zcl 0.1 released

Just wanted to get it out of the way so I did a bit of house-keeping on my isp web page and created a zcl home page. It pretty much explains everything that could be said here.

As part of the house-keeping I created a basic but functional software index page to consolidate the mess of stuff I had there and moved all the download files to a common location.

Update: I went looking for that option that turns on warnings for functions without prototypes ... and found a bug. enqueueReadImage() can't work in 0.1.

Anyway I found something to do: get it going on parallella.

toArray or not to array

So I guess I found some more stuff to play around with. It's a really lovely day here but I know town is going to be packed a bit more than I'd like and I've kinda had enough of going out for the moment. I guess i'll see after lunch on that.

First I thought i'd fix the array getter methods. The OpenCL get methods are built in a way where if you don't know the size of a field you have to call it twice - once to get the size and the next to get the content. There didn't seem to be much point exposing this to the Java side of things as I had initially done and I went with just doing that directly in the C code and returning a newly allocated and correctly sized result. My original thoughts were that perhaps buffers could be re-used for multiple gets but in reality it just isn't that useful: for small buffers it doesn't matter and for large ones you need to find out the size anyway and any management overheads (thread-specific etc, and simply having memory sitting around doing nothing) of reusable buffers is going to swamp allocation and GC.

I also realised I could fix one of the last bastions of the exposed native pointers and change the long getInfoP methods to return an object directly to:

 native public <T extends CLObject> T getInfoP(int param, Class<T> klass);

Which is kind of nice. Actually getInfoP() was hidden by type-specic getters but doing it this way (and particularly for the array types) saved even more code in the Java side for a minimal cost on the C side (actually I ended up saving code by reorganising the array getters).

Then I thought about whether I could add native array types to the CLCommandQueue interfaces. e.g.

    native public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
        long mem_offset, long size,
        byte[] buffer, long buf_offset,
        CLEventList wait,
        CLEventList event) throws CLException;

In addition to the interface that uses nio buffers.

The tricky bit is that these can run asynchronously so you can't use the GetPrimitiveArrayCritical() calls and you're basically left with either manually copying them using Get*ArrayRegion() or using Get*ArrayElements() which just seems to copy them on hotspot anyway.

As an experiment I tried the latter. Actually it ends up copying both ways which is a bit of a waste.

When called without blocking I use an event callback to await completion and then release the array back to Java. Strictly speaking I should also do the same for the Buffer versions so that the Buffer doesn't get GC'd while it's running but that's something I think can be left to the programmer to keep track of.

I tried a test program which just did many calls followed by a flush each time and actually performance wasn't too bad relative to the Buffer version. Maybe 10-20% slower (which is ok since accessing arrays is faster and simpler than Buffers in java). But then I tried a silly example of moving the flush outside of the loop. Ok, now it's 4x slower and god knows how much memory it ends up swallowing whilst executing.

So I followed up by trying the GetArrayRegion interface. This is a little bit faster but nothing to write home about.

At this point I think i'll just keep the binding and api smaller and leave it with using a ByteBuffer (sigh, which i still need to fix the endianess of) but i'll save the code for maybe later.

Actually probably the most surprising thing is just how slow the OpenCL stuff is here. This is only using the CPU driver so there's no weird memory busses to go over (even if this wasn't an apu). It's about 100x slower than copying a ByteBuffer to a byte array the same number of times. I thought it might be because the calls are non-blocking, but making them blocking only makes it worse. I tested the JNI overhead too by simply nooping out the clEnqueueReadBuffer call on the array Region version and that is only about 2x slower than ByteBuffer.get().

Yeah ...

ZCL == feature complete

I managed to completely fill out the OpenCL binding i've been playing with: wrote accessors for all the properties which made sense (only the refcount stuff omitted, CLDevice has a fuckload), cleaned up a few bits pieces, added lots of validity checks, fixed up some portability things, added some (incomplete) javadocs, license headers, a README, and a couple of incomplete hacked-up makefiles which build all 4 outputs (jar, javadoc, jni library, source distribution) in under 5 seconds.

Counting semi-colons it's about 1KLOC of Java and 1K3LOC of C; which seems quite reasonable for 100% api coverage, and KLOC was a big part of the experiment I was conducting. And even then about 400 of those Java lines are just a copy of the defines from cl.h. Kinda lost-track of how much time i've spent on it at this point - something like 4 not-quite-full days work.

But now i'm kinda bored with that toy.

Don't really feel like testing enough of it to get it to bee-ta state at this point. And that's the point I guess; there isn't a point to any of it. The end-goal wasn't the interesting bit.

Journey's over, what next?

But I spose ... i'll dump it somewhere when I feel like spitting out a home-page for it.

Native kernels too?

I kept poking from the previous post and ended up getting native kernels going as well. I'm not really sure how useful they are but it's nice to come up with a neat solution.

It took me a while to grok the interface to clEnqueueNativeKernel but it seems to make sense.

This is the result I managed:

  public interface CLNativeKernel {
    public void invoke(Object[] args);
  }

  class CLCommandQueue {
      public native void enqueueNativeKernel(
          CLNativeKernel kernel,
          CLEventList waiters,
          CLEventList events,
          Object... args) throws CLException;
  }

Which leads to a relatively clean usage:

 CLBuffer mem = cl.createBuffer(0, 1024 * 4, null);

 q.enqueueNativeKernel((Object[] args) -> {
    System.out.printf("native kernel invoked %s\n", Thread.currentThread());
    for (Object o : args) {
        System.out.printf(" %s = %s\n", o.getClass().getName(), o);
    }
  }, null, null, mem, 10, mem, 10L);

Produces:

native kernel invoked Thread[Thread-0,5,main]
 java.nio.DirectByteBuffer = java.nio.DirectByteBuffer[pos=0 lim=4096 cap=4096]
 java.lang.Integer = 10
 java.nio.DirectByteBuffer = java.nio.DirectByteBuffer[pos=0 lim=4096 cap=4096]
 java.lang.Long = 10

The tricky bit is getting the memory handled. clEnqueueNativeKernel takes cl_mem arguments as input but then remaps them to physical (virtual) memory pointers when invoking the kernel. The only equivalent of a pointer in Java is a ByteBuffer ... but that also needs a length.

But basically I just copy over the jobject references from the jobject array and change any CLMemory classes to be the cl_mem they point to. In the native kernel hook I then have to remap the provided pointers of any CLMemory instances to direct ByteBuffers, and I obtain the actual memory size using clGetMemObjectInfo(). Because the native kernel hook can only take one set of arguments I fudge it by internally using argument 0 as a structure which contains all the copies of stuff I need and then free it afterwards. It does force the java code deal with some of the bytebuffer details but the only alternatives I can think of get pretty messy and actually doing lots of processing on memory buffers isn't something you should be doing from any native kernel to start with. They only work on CPU targets (APU?) anyway.

I did hit an issue in that AttachCurrentThread() was attaching to another native thread this time; so I tried using AttachCurrentThreadAsDaemon() instead. That may actually not be a good idea but it depends on whether a given OpenCL implementation is using thread pools or not. I guess?

Anyway, i'm fairly pleased with the result here.

on JNI callback reference handling

After the previous post detailing some issues with handling callback reference handling I had another look at it this evening.

First for the clBuildProgram() function I just deleted the global reference in the callback. I tried to identify reference leaks using the netbeans memory profiler but it was a little difficult to interpret the results. For starters running the demo routine in a loop didn't result in loop-number reference leaks as one would expect (or even loop-number of reference creations oddly enough; may be related to hotspot and/or it being a static method) ... anyway I think it should work regardless except in the specific case where OpenCL doesn't actually call the notify callback for whatever reason: it is unclear from the specification if it MUST always call it for example. I'm just going to have to assume if that ever happens the system is in such a state then adding a leak is of no practical importance.

Then took at look at the clCreateContext() issue which seemed a bit trickier. On a hunch I looked up how weak references work from JNI and at first I didn't see anything useful but whilst poking around a tidy solution became apparent.

All I have to do is save the original reference to any notify function in the CLContext on the Java side. This lets Java handle the reference as it normally would and any notify object should automatically have the same lifetime as the reference to CLContext.

From the (rather badly formatted) JNI document:

Weak Global References

Weak global references are a special kind of global reference. Unlike normal global references, a weak global reference allows the underlying Java object to be garbage collected. Weak global references may be used in any situation where global or local references are used. When the garbage collector runs, it frees the underlying object if the object is only referred to by weak references. A weak global reference pointing to a freed object is functionally equivalent to NULL. Programmers can detect whether a weak global reference points to a freed object by using IsSameObject to compare the weak reference against NULL.

Weak global references in JNI are a simplified version of the Java Weak References, available as part of the Java 2 Platform API ( java.lang.ref package and its classes).

Clarification    (added June 2001)

Since garbage collection may occur while native methods are running, objects referred to by weak global references can be freed at any time. While weak global references can be used where global references are used, it is generally inappropriate to do so, as they may become functionally equivalent to NULL without notice.

While IsSameObject can be used to determine whether a weak global reference refers to a freed object, it does not prevent the object from being freed immediately thereafter. Consequently, programmers may not rely on this check to determine whether a weak global reference may used (as a non- NULL reference) in any future JNI function call.

To overcome this inherent limitation, it is recommended that a standard (strong) local or global reference to the same object be acquired using the JNI functions NewLocalRef or NewGlobalRef , and that this strong reference be used to access the intended object. These functions will return NULL if the object has been freed, and otherwise will return a strong reference (which will prevent the object from being freed). The new reference should be explicitly deleted when immediate access to the object is no longer required, allowing the object to be freed.

  public class CLContext extends CLObject {

    final CLContextNotify notify;

    CLContext(long p, CLContextNotify notify) {
      super(p);
      this.notify = notify;
    }
  }

  -  data = (*env)->NewGlobalRef(env, jnotify);
  +  data = (*env)->NewWeakGlobalRef(env, jnotify);

  +  jnotify = (*env)->NewLocalRef(env, jnotify);
  +  if (!jnotify) {
  +    fprintf(stderr, "cl_context notify called after object death\n");
  +    return;
  +  }

OpenCL binding - a bit more tweaking

Had a bit of another look at the OpenCL binding I was working on. I wasn't happy that some of the public interfaces still uses long[] arrays to represent intptr_t arrays - specially for property lists. So I made a bit more java-ish. It's still a bit clumsy but it's about as good as it's going to get.

    public static native CLContext createContext(long[] properties,
            CLDevice[] devices,
            CLContextNotify notify) throws CLRuntimeException;

Becomes:

    public static native CLContext createContext(CLContextProperty[] properties,
            CLDevice[] devices,
            CLContextNotify notify) throws CLRuntimeException;

Properties all inherit from a base class:

public class CLProperty {
    protected final long tag;
    protected final long value;

    protected CLProperty(long tag, long value) {
        this.tag = tag;
        this.value = value;
    }
}

This is so the JNI code only needs to deal with one type of object. Then I have factory methods for the various property types.

public class CLContextProperty extends CLProperty {

    ...

    public static CLContextProperty CL_CONTEXT_PLATFORM(CLPlatform platform) {
        return new CLContextProperty(CL.CL_CONTEXT_PLATFORM, platform.p);
    }

    ...

}

Although I might make the names more java-friendly.

So based on the createContext interfaces above, one changes:

  cl = createContext(new long[] { CL.CL_CONTEXT_PLATFORM, platform.p, 0 },
                     new CLDevice[] { dev },
                     null);

to:

  cl = createContext(new CLContextProperty[] { CLContextProperty.CL_CONTEXT_PLATFORM(platform) },
                     new CLDevice[] { dev },
                     null);

It's not like it saves typing but it is type-safe, and you don't have to remember to put the closing 0 tag on the end of the list. Perhaps the factory methods should sit on CLContext for that matter.

Callbacks, Leaks, Lambdas

Another part I looked into implementing was the callback methods from C to Java, such as the one passed to createContext or buildProgram.

This is mostly straightforward - just pass a hook function to the OpenCL call which locates an environment and invokes the callback function on an interface. There is no need to support a 'user data' field for the java side, so that is just used to pass a global reference to the interface itself.

If one considers the generic interface used for build callbacks:

public interface CLNotify<T< {
    public void notify(T source);
}

The C hook is relatively straightforward ...

static void build_notify_hook(cl_program prog, void *data) {
    jobject jnotify = data;
    jobject source;
    JNIEnv *env;
    jlong lprog = (jlong)prog;

    if ((*vm)->GetEnv(vm, (void *)&env, JNI_VERSION_1_4) != 0
        && (*vm)->AttachCurrentThread(vm, (void *)&env, NULL) != 0) {
        fprintf(stderr, "Unable to attach java environment\n");
        return;
    }

    source = (*env)->NewObjectA(env, classid[PROGRAM], new_p[PROGRAM], (void *)&lprog);
    if (!source)
        return;

    (*env)->CallVoidMethodA(env, jnotify, CLNotify_notify, (void *)&source);
}

(FIXME: this may need to detach the thread also). (FIXME: this may need to de-ref jnotify)

One notices that the callback simply creates a new CLProgram object instance to the pass the pointer to Java. This means that OpenCL handles may map to more than one Java object: this goes some way to validating my decision to stick with simple holder objects rather than trying to keep some data copied to the Java side. Although it wouldn't be that difficult to track object instances if necessary: instead of calling NewObject() invoke a factory method which handles the object instances. Albeit at the cost of duplicating the reference tree in Java.

Another bonus i didn't realise is that the way lambdas are implemented allows these to be used from the Java side without the JNI needing to know anything about it. I think I did read about this at some point but it's been a while and I forgot about it. I had a look at a dissassemby of the class file and it's just using invokedymanic to create an interface object which is just a function pointer rather than having to create an instance of an abstract class.

So e.g. this works:

  prog.buildProgram(new CLDevice[]{dev}, null,
    (CLProgram source) -> {
        System.out.printf("Build notify, status = %d\nlog:\n",
            source.getBuildInfoInt(dev, CL_PROGRAM_BUILD_STATUS));
        System.out.println(source.getBuildInfoString(dev, CL_PROGRAM_BUILD_LOG));
    });

The one very big caveat for all of the above ... is that I haven't worked out a clean way to avoid leaking the notify instance object. This is because the OpenCL api specifies that these callback functions may be invoked asynchronously and/or from other threads.

Thinking aloud:

For the specific case of clBuildProgram and friends it looks like the notify function is only ever (and always) called once and I can thus deref the interface in the hook routine. If I pass both the CLProgram object and CLNotify interface to the hook routine I can keep the CLProgram instance unique anyway ... (And to be honest i'm not sure how useful this mechanism is to start with since it's easier just to compile synchronously and check the return code / exception).

But CLContext has it's own notify function too which needs to live as long as the CLContext so I can't use the same trick there. At first I thought of creating an set/remove listener interface that just keyed everything off the point value and tracking the listeners in Java. But that doesn't work because presumably it's possible to get a callback call without ever getting a context. I guess I could use the listener itself as a key and provide a static native clearNotify() method which must be called explicitly but it gets a bit messy for a few reasons.

struct notify_info {
  int id;
  jobject jnotify;
};

clCreateContext(..., jobject jnotify) {
...
  lock {
   info = malloc();
   info.id = getsequence();
   info.jnotify = NewGlobalRef(jnotify);
   listeners.add(info);
  }
...
  clCreateContext(..., create_context_hook, (void *)id);
...
}

create_context_hook(..., void *data) {
  int id = (int)data;

  lock {
     info = listeners.find(id);
     if (info) {
        ... invoke info.jnotify;
     }
  }
}

clear_context_notify(..., jobject jnotify) {
   lock {
      info = listeners.find(jnotify);
      if (info) {
         deleteGlobalRef(info.jnotify);
         listeners.remove(info);
      }
   }
}

Yeah, messy. A bunch of it could be (synchronous) static Java methods, but it just isn't particularly elegant either way.

Again i'm not sure how useful implementing this precise interface is anyway: it may just as well do to implement a completely separate system which funnels all events through a global event handler mechanism.

javafx + internet radio = partial success?

After a long and pointless goose-chase with different versions of ffmpeg and libav ... I got my internet radio thing to work from JavaFX.

It turned out to be at least partly a problem with my proxy code. The ice-cast stream doesn't include a Content-Length header (because you know, it's a stream), so this was causing libfxplugins to crash as in my previous post on the subject. Shouldn't really cause a crash - at most it should just throw a protocol error exception? Not sure if it has javafx security implications (e.g. if media on a web page through the web component goes through the same mechanism it certainly does - Update: I filed a bug in jira, and it seems to have been escalated - even at this late stage of java 8).

If I add a made-up content-length value at least I get the music playing, .. but the 'user experience' isn't very good because it seems to change how much it pre-buffers depending on the reported length (rather strange i think). So if i report 1-10M it starts after a few seconds but if i report some gigantic number it buffers for minutes (or maybe it never starts, i lost patience). The problem is that as the data length is hounoured (as it should be) which means that too small a number causes the stream to finish quickly. And it seems to read the same stream twice at the same time which is also odd.

To be blunt I find it pretty perplexing 'in this day and age' that javafx doesn't support streaming media to start with. Or that g-"so-called-streamer" seems to be the reason for this. If I could work out how to compile libfxplugins.so from the openjfx dist I would have a poke, but alas ...

I was going to go ahead with the application anyway but I think because of this streaming issue I wont - or if I do I wont be using javafx to do the decoding which is a bit of a pain.

Update 2: The bug has just been patched. 4 working days from reporting and if the openjdk schedule is still correct, just 2 weeks from release. So yeah I guess it was reasonably important and even just mucking about has value.

Simple Java binding for OpenCL 1.2

Well I have it to the point of working - it still needs some functions filled out plus helper functions and a bit of tweaking but it's mostly there. So far under 2KLOC of C and less of Java. I went with the 'every pointer is 64-bits' implementation, using non-static methods, and passing objects around to the JNI rather than the pointers (except for a couple of apis). This allows me to implement the raw interface fully in C with just an 'interface' in Java - and thus write a lot less code.

Currently i'm mapping a bit closer to the C api than JOCL does. I'm using only using ByteBuffers to transfer memory asynchronously, for any other array arguments i'm just using arrays.

This example is with the raw api with no helpers coming in to play - there are some obvious simple ones to add which will make it a bit more comfortable to use.

// For all the CL_* constants
import static au.notzed.zcl.CL.*;

...

  CLPlatform[] platforms = CLPlatform.getPlatforms();
  CLPlatform plat = platforms[0];
  CLDevice dev = plat.getDevices(CL_DEVICE_TYPE_CPU)[0];
  CLContext cl = plat.createContext(dev);
  CLCommandQueue q = cl.createCommandQueue(dev, 0);

  CLBuffer mem = cl.createBuffer(0, 1024 * 4, null);

  CLProgram prog = cl.createProgramWithSource(
    new String[] {
      "kernel void testa(global int *buffer, int4 n, float f) {" +
      " buffer[get_global_id(0)] = n.s1 + get_global_id(0);" +
      "}"
    });

  pog.buildProgram(new CLDevice[]{dev}, null, null);

  CLKernel k = prog.createKernel("testa");

  ByteBuffer buffer = ByteBuffer.allocateDirect(1024 * 4).order(ByteOrder.nativeOrder());
   
  k.set(0, mem);
  k.set(1, 12, 13, 14, 15);
  k.set(2, 1.3f);
   
  q.enqueueWriteBuffer(mem, CL_FALSE, 0, 1024 * 4, buffer, 0, null, null);
  q.enqueueNDRangeKernel(k, 1, new long[] { 0 }, new long[] { 16 }, new long[] { 1 }, null, null);
  q.enqueueReadBuffer(mem, CL_TRUE, 0, 1024 * 4, buffer, 0, null, null);
  q.finish();

  IntBuffer ib = buffer.asIntBuffer();
   
  for (int i=0;i<32;i++) {
    System.out.printf(" %3d = %3d\n", i, ib.get());
  }

Currently CLBuffer (and CLImage) is just a handle to the cl_mem - it has no local meta-data or a fixed Buffer backing. The way JOCL handles this is reasonably convenient but i'm still yet to decide whether I will do something similar. Whilst it may be handy to have local copies of data like 'width' and 'format', I'm inclined to just have accessors which invoke the GetImageInfo call instead - it might be a bit more expensive but redundant copies of data isn't free either.

I'm not really all that fond of the way JOCL honours the position() of Buffers - it kind of seems useful but usually it's just a pita. And manipulating that from C is also a pain. So at the moment I treat them as one would treat malloc() although I allow an offset to be used where appropriate.

Such as ...

public class CLCommandQueue {
  ...
   native public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
      long mem_offset, long size,
      Buffer buffer, long buf_offset,
      CLEventList wait,
      CLEventList event) throws CLException;
  ...
}

Compare to the C api:

extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueWriteBuffer(cl_command_queue   /* command_queue */, 
                     cl_mem             /* buffer */, 
                     cl_bool            /* blocking_write */, 
                     size_t             /* offset */, 
                     size_t             /* size */, 
                     const void *       /* ptr */, 
                     cl_uint            /* num_events_in_wait_list */, 
                     const cl_event *   /* event_wait_list */, 
                     cl_event *         /* event */) CL_API_SUFFIX__VERSION_1_0;

In C "ptr" can just be adjusted before you use it but in Java I need to pass buf_offset to allow the same flexibility. It would have been nice to be able to pass array types here too ... but then I realised that these can run asynchronous which doesn't work from jni (or doesn't work well).

I'm still not sure if the query interface is based only on the type-specific queries implemented in C or whether I have helpers for every value on the objects themselves. The latter makes the code size and maintenance burden a lot bigger for questionable benefit. Maybe just do it for the more useful types.

Haven't yet done the callback stuff or native kernels (i don't quite understand those yet) but most of that is fairly easy apart from some resource tracking issues that come in to play.

Of course now i've done 90% of the work i'm not sure i can be fagged to do the last 10% ...