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% ...

more on JNI overheads

I wrote most of the OpenCL binding yesterday but now i'm mucking about with simplifying it.

I've experimented with a couple of binding mechanisms but they have various drawbacks. They all work in basically the same way in that there is an abstract base class of each type then a concrete platform-specific implementation that defines the pointer holder.

The difference is how the jni C code gets hold of that pointer:

Passed directly

The abstract base class defines all the methods, which are implemented in the concrete class, which just invokes the native methods. The native methods may be static or non-static.

This requires a lot of boilerplate in the java code, but the C code can just use a simple cast to access the CL resources.

C code performs a field lookup

The base class can define the methods directly as native. The concrete class primarily is just a holder for the pointer value.

This requires only minimal boiler-plate but the resources must be looked up via a field reference. The field reference is dependent on the type though.

C code performs a virtual method invocation.

The base class can define the methods directly as native. The concrete class primarily is just a holder for the pointer value.

This requires only minimal boiler-plate but the resources must be looked up via a method invocation. But here the field reference is independent on the type.

The last is kind of the nicest - in the C code it's the same amount of effort (coding wise) as the second but allows for some polymorphism. The first is the least attractive as it requires a lot of boilerplate - 3 simple functions rather than just one empty one.

But, a big chunk of the OpenCL API is dealing with mundane things like 'get*Info()' lookups and to simplify it's use I came up with a number of type-specific calls. However rather than write these for every possible type I pass a type-id to the JNI code so a single function works. This works fine except that I would like to have a separate CLBuffer and CLImage object - and in this case the second implementation falls down.

To gain more information on the trade-off involved I did some timing on a basic function:

  public CLDevice[] getDevices(long type) throws CLException;

This invokes clGetDeviceIDs twice (first to get the list size) and then returns an array of instantiated wrappers for the pointers. I invoked this 10M times for various binding mechanisms.

Method                   Time
 pass long                13.777s
 pass long static         14.212s
 field lookup             14.060s
 method lookup            16.252s

So interesting points here. First is that static method invocations appear to be slower than non-static even when the pointer isn't being used. This is somewhat surprising as 'static' methods seem to be quite popular as a mechanism for JNI binding.

Second is that a field lookup from C isn't that much cost compared to a field lookup in Java.

Lastly, as expected the method lookup is more expensive and if one considers that the task does somewhat more than the pointer resolution then it is quite significantly more expensive. So much so that it probably isn't the ideal solution.

So ... it looks like I may end up going with the same solution I've used before. That is, just use the simple field lookup from C. Although it's slightly slower than the first mechanism it is just a lot less work for me without a code generator and produces much smaller classes either way. I'll just have to work out a way to implement the polymorphic getInfo methods some other way: using IsInstanceOf() or just using CLMemory for all memory types. In general performance is not an issue here anyway.

I suppose to do it properly I would need to profile the same stuff on 32-bit platforms and/or android as well. But right now i don't particularly care and don't have any capable hardware anyway (apart from the parallella). I wasn't even bothering to implement the 32-bit backend so far anyway.

Examples

This is just more detail on how the bindings work. In each case objects are instantiated from the C code - so the java doesn't need to know anything about the platform (and is thus, automatically platform agnostic).

First is passing the pointer directly. Drawback is all the bulky boilerplate - it looks less severe here as there is only a single method.

public abstract class CLPlatform extends CLObject {

    abstract public CLDevice[] getDevices(long type) throws CLException;

    class CLPlatform64 extends CLPlatform {
        final long p;

        CLPlatform64(long p) {
            this.p = p;
        }

        public CLDevice[] getDevices(long type) throws CLException {
            return getDevices(p, type);
        }

        native CLDevice[] getDevices(long p, long type) throws CLException;
    }

    class CLPlatform32 extends CLPlatform {
        final int p;

        CLPlatform32(int p) {
            this.p = p;
        }

        public CLDevice[] getDevices(long type) throws CLException {
            return getDevices(p, type);
        }

        native CLDevice[] getDevices(int p, long type) throws CLException;
    }
}

Then having the C lookup the field. Drawback is each concrete class must be handled separately.

public abstract class CLPlatform extends CLObject {
    native public CLDevice[] getDevices(long type) throws CLException;

    class CLPlatform64 extends CLPlatform {
        final long p;

        CLPlatform64(long p) {
            this.p = p;
        }
    }

    class CLPlatform64 extends CLPlatform {
        final long p;

        CLPlatform64(long p) {
            this.p = p;
        }
    }

    class CLPlatform32 extends CLPlatform {
        final int p;

        CLPlatform32(long p) {
            this.p = p;
        }
    }
}

And lastly having a pointer retrieval method. This has lots of nice coding benefits ... but too much in the way of overheads.

public abstract class CLPlatform extends CLObject {
    native public CLDevice[] getDevices(long type) throws CLException;

    class CLPlatform64 extends CLPlatform implements CLNative64 {
        final long p;

        CLPlatform64(long p) {
            this.p = p;
        }

        long getPointer() {
            return p;
        }
    }

    class CLPlatform64 extends CLPlatform implements CLNative32 {
        final int p;

        CLPlatform64(int p) {
            this.p = p;
        }

        int getPointer() {
            return p;
        }
    }
}

Or ... I could of course just use a long for storage on 32-bit platforms and be done with it - the extra memory overhead is pretty much insignificant in the grand scheme of things. It might require some extra work on the C side when dealing with a couple of the interfaces but it is pretty minor.

With that mechanism the worst-case becomes:

public abstract class CLPlatform extends CLObject {
    final long p;

    CLPlatform(long p) {
        this.p = p;
    }

    public CLDevice[] getDevices(long type) throws CLException {
        return getDevices(p, type);
    }

    native CLDevice[] getDevices(long p, long type) throws CLException;
}

Actually I can move 'p' to the base class then which simplifies any polymorphism too.

I still like the second approach somewhat for a hand-coded binding since it keeps the type information and allows all the details to be hidden in the C code where it is easier to hide using macros and so on. And the java becomes very simple:

public abstract class CLPlatform extends CLObject {
    CLPlatform(long p) {
        super(p);
    }

    public native CLDevice[] getDevices(long type) throws CLException;
}

CLEventList

Another problematic part of the OpenCL api is cl_event. It's actually a bit of a pain to work with even in C but the idea doesn't really map well to java at all.

I think I came up with a workable solution that hides all the details without too much overheads. My initial solution is to have a growable list of items (the same as JOCL) that was managed on the Java side. It's a bit messy on the C side but really messy on the Java side:

public class CLEventList {
   static class CLEventList64 {
      int index;
      long[] events;
   }
}

...
   enqueueSomething(..., CLEventList wait, CLEventList event) {
       CLEventList64 wait64 = (CLEventList64)wait;
       CLEventList64 event64 = (CLEventList64)event;

       enqueueSomething(...,
           wait64 == null ? 0 : wait64.index, wait64 == null ? null : wait64.events,
           event64 == null ? 0 : event64.index, event64 == null ? null : event64.events);

       if (event64 != null) {
           event64.index+=1;
       }
   }

Yeah, maybe not - for the 20 odd enqueue functions in the API.

So I moved most of the logic to the C code - actually the logic isn't really any different on the C side it just has to do a couple of field lookups rather than take arguments, and I added a method to record the output event.

public class CLEventList {
   static class CLEventList64 {
      int index;
      long[] events;
      void addEvent(long e) {
        events[index++];
      }
   }
}

...
   enqueueSomething(..., CLEventList wait, CLEventList event) {
       enqueueSomething(...,
           wait,
           event);
   }

UserEvents are still a bit of a pain to fit in with this but I think I can work those out. The difficulty is with the reference counting.