Showing posts with label opencl. Show all posts
Showing posts with label opencl. Show all posts

Saturday, 9 September 2017

zcl 0.6

Yes it still lives. I've just uploaded an update to zcl

A bunch of bugfixes, new build system, more robustness, and OpenCL 2.1 support.

There are still some thing i'm experimenting with - primarily the functional/task stuff as it's just not flexible enough - but it's stable and robust and easy to work with so i'm no longer using JOCL for anything at work.

On a personal note I still haven't really gotten back into hacking and i had a short sojourn into facebookland so i haven't had much to write about. It's mostly been work, very poor sleep, and drinking! Oh and I started wearing kilts ...

Wednesday, 17 February 2016

vulkan is out at last

I'm just starting to think about hacking again so that delay might not have been so bad for me. I haven't looked at it yet though.

Oh, and what a nice birthday present too!

Update: Not that i've had time to look but still no AMD drivers for linux; although if this means an end to fglrx then that will be worth it.

Tuesday, 1 December 2015

more termz

I tried a few more variants on the OpenCL rendering code but none are all that fast - the overheads kill it and whilst it does free the CPU up a little bit it isn't much. Probably not worth more time unless I look at OpenGL instead.

I added resizing. It meant I had to add some locking to the snapshot routine but it only needs to lock around the resize operation so it adds almost no overhead to normal operation (merely detecting a resize has occurred). I'm not yet sure what's supposed to happen with saved cursors and alternate screens when a resize occurs, probably just clip to size. Unfortunately there seems to be no way to set the WM_NORMAL_HINTS on the javafx stage, so there's no way to make it size to cells properly.

I did a bunch of benchmarking and profiling. One thing I tried was another test to compare to xterm - now at full-screen. Running "find . -name '*.c' | xargs cat" from the root of the linux 3.19.8 source tree. After a couple of runs this is about 25 seconds on termz, and 16 minutes in xterm. Well. Yeah it's a silly test but all those ls -l's during the day add up.

Looking at the memory profiler it doesn't really use too much heap during operation, just a few MB and most of that is the image and javfx. I mean nor should it, there isn't data structures to maintain. But having multiple compilers in ram (jvm, OpenCL), their generated output, and all the added overhead of the runtime support needed for those really adds up so it's very very fat in practice. I guess if multiple terms ran on the same jvm it would be ok.

enough yet?

So at this point i'm not really sure what i'll do with it. I'll probably poke at the edges when i'm bored and eventually when I get around to it I will dump what I have to my software site as the result of a weekend-and-a-bit-hack.

After that I'm not sure. It's actually quite functional and robust already (well, compared to the effort in) and wouldn't take much more work to turn it into a usable terminal for me - add some scrollback (pretty simple), mouse selection stuff (not that hard), and sort out some of the keyboard details (reading obtuse documents and testing). So maybe that will happen.

If I got that far adding the "10x20" typeface would probably be on the cards. Fixed-size outline fonts would be possible by just pre-rendering them but to me they just aren't terminal fonts.

Anything further such as a re-usable term component (which might actually be of use to the world) would require substantially more work on the i18n side of things and I don't feel like learning enough to do that properly.

Make

I did a bit poking at the java makefile stuff and it's to the point where i'm using it for out-of-ide builds of termz and zcl and will look into using it on other projects. That's the best way to find bugs/what works and what doesn't and previous attempts never got that far. For all it's gnumakefile obtuseness it's really rather compact at under 200 lines excluding comments and I didn't put any effort into making it particularly small. And that includes targets for javadocs, source jars, and binary builds.

Sunday, 29 November 2015

opencl termz

I was just going to "try something out" while I waited for the washing to finish ...

... so after a long and full day of hacking ...

The screenshot is from an OpenCL renderer. Each work item processes one output pixel and adds any attributes on the fly, somewhat similar in effect to how hardcoded hardware might have done it. I implemented a 'fancy' underline that leaves a space around descenders. The font is a 16x16 glyph texture of iso-8859-1 characters. I haven't implemented colour but there's room for 16.

On this kaveri machine with only one DIMM (== miserable memory bandwidth) the OpenCL routine renders this buffer in about 35-40uS. This doesn't sound too bad but it takes 3uS to "upload" the cell table input, and 60uS to "download" the raster output (and this is an indexed-mode 8-bit rather than RGBA which is ~2x slower again), but somehow by the time it's all gone through OpenCL that's grown to 300-500uS from first enqueue to final dequeue. Then add on writing to JavaFX (which converts it to BGRA) and it ends up ~1200uS.

I'm using some synchronous transfers and just using buffer read/write so there could be some improvements but the vast majority of the overheads are due to the toolkit.

So I guess that's "a bit crap" but it would still be "fast enough". For comparison a basic java renderer that only implements inverse is about 1.5x slower overall.

But for whatever reason the app still uses ~8% cpu even when not doing anything; and that definitely isn't ok at all. I couldn't identify the cause. Another toolkit looks like a necessity if it ever went beyond play-thing-toy.

I got bored doing the escape codes around "^[ [ ? Ps p" so it's broken aplenty beyond the bits I simply implemented incorrectly. But it's only a couple days' poking and just 1K3LOC. While there is ample documentation on the codes some of the important detail is lacking and since i'm not looking at any other implementation (not even zvt) i have to try/test/compare bits to xterm and/or remember the fiddly bits from 15 years ago (like the way the cursor wrapping is handled). I also have most of the slave process setup sorted beyond just the pty i/o - session leaders, controlling terminals, signal masks and signal actions, the environment. It might not be correct but I think all the scaffolding is now in place (albeit only for Linux).

FWIW a test i've been using is "time find ~/src" to measure elapsed time on my system - after a couple of runs to load everything into the buffer cache this is a consistent test with a lot of spew. If I run it in an xterm of the same size this takes ~25s to execute and grinds big parts of the desktop to a near halt while it's active. It really is abysmal behaviour given the modern hardware it's on (however "underpowered" it's supposed to be; and it's considerably worse on a much much faster machine). The same test in 'termz' takes about 4.5s and you'd barely know it was running. Adding a scrollback buffer would increase this (well probably, and not by much) however this goes through a fairly complete UTF-8 code-path otherwise.

The renderer has no effect on the runtime as it is polled on another thread (in this instance via the animation pulse on javafx). I don't use locks but rely on 'eventual consistency'. Some details must be taken atomically as part of the snapshot and these are handled appropriately.

Right now I feel like i've had my fill for now with this. I'm kinda interested, but i'm not sure if i'm interested enough to finish it sufficiently to use it - like pretty much all my hacking hacked up hacks. Time will be the teller.

Friday, 27 November 2015

termz

Every now and then I think about the sad state of affairs regarding terminal emulators for X11. It's been a bit of a thing for a while - it's how i ended up working at Ximian.

I stopped using gnome-terminal when i stopped working on it and went back to xterm. I never liked rxvt or their ilk and all of the 'desktop environment' terminal emulators are pretty naff for whatever reason.

xterm works and is reliable but with recent (being last 10 years) X Windows System servers the text rendering performance plummeted and even installing the only usable typefaces (misc-fixed 6x13, and 10x20, and sometimes xterm itself) became a manual job. Whilst performance isn't bad on this kaveri box I also use an uber-intel machine with a HD7970 where both emacs and xterm runs like an absolute pig whenever any GL applications are running, and it isn't even very fast otherwise (i'm talking whole desktop grinding to a halt as it redraws exposes at about 1 character column per SECOND). It's an "older" distribution so that may have something to do with it but there is no direct indication why it's so horrible (well apart from the AMD driver but i have no choice for that since it's used for OpenCL dev). I might upgrade it to slackware next year.

Ho hum.

Anyway I started poking last night at a basic xterm knockoff and got to the point of less sort of running inside it and now i'm thinking about ways I might be able to implement something a bit more complete. I'm working in Java and have a tiny bit of JNI to get the process going and handle some ioctl stuff (which seems somewhat easier now than it was in zvt, but portability is not on the agenda here).


TermZ? Glyphs are greymaps extracted directly from the PCF font.

zvt

When I wrote ZVT the primary goal was performance and to that end considerable effort was expended on making a terminal state machine which implemented zero-copy and zero-garbage algorithms. zero-copy is always a good thing but the zero-garbage was driven by the very slow malloc on Solaris at the time and my experience with Amiga memory management.

Another part of the puzzle was display and the main mechanism was inspired by some Amiga terminal emulators that used COPPER lists to re-render rows to the screen in arbitrary order without requiring them to be re-ordered in memory (memory bandwidth was a massive bottleneck when using pre 1985-spec hardware in 199x). I used a cyclic double-linked (exec) list of rows and to implement a scroll I just moved a row from the start to the end of the list which takes 8 pointer updates and a memset to clear it (and it also works for partial screen scrolls). By tracking the last row a given one was actually displayed at I could -at-any-point-later- attempt to create an optimal screen-update sequence including using blits for scrolling and minimising glyph redraws to only those that had changed. The algorithm for this was cheap and reliable if a little fiddly to get correct.

This last point is important as it allows the state machine to outpace the screen refresh rate which always becomes the largest bottleneck for terminal emulators in 'toolkit' environments. This is where it got all it's performance from.

new hardware, new approach

Thinking about the problem with current hardware my initial ideas are a little bit different.

I still quite like the linked list storage for the state machine and may go back to it but my current idea is instead to store a full cell-grid for the displayable area. I can still make full-screen scrolling just as cheap using a simple cyclic row trick (infact, even cheaper) but sub-region scrolling would require memory copies - but at the resolution of 4-bytes-per-glyph these are insanely cheap nowadays.

This is the most complex part of the emulator since it needs to implement all the control codes and whatnot - but for the most part thats just a mechanical process of implementing enough of them to have something functional.

I would also approach rendering from an entirely different angle. Rather than go smart i'm going wide and brute-forcing a simpler problem. At any given time - which can be throttled based on arbitrary metrics - I can take a snapshot of the current emulator screen and then asynchronously convert that to a screen display while the emulator continues to run on it's own thread.

For a basic CPU renderer it may still require some update optimisation but given it will just be trivial cell fonts to copy it probably wont be appreciably cheaper to scroll compared to just pasting new characters every time. And obviously this is utterly trivial code to implement.

The ultimate goal (and why the fixed-array grid backing is desirable) would be to use OpenCL or OpenGL (or more likely Vulkan if it ever gets here) to implement the rendering as a single pass operation which touches each output pixel only once. This would just take the raw cell-sized rectangle of the terminal state machine as it's only variable input and produce a fully rendered and styled framebuffer as the result. Basically just render the cells as a low-res nearest-neighbour texture lookup into a texture holding the glyphs. The former is a tiny tiny texture in GPU terms and rendering a single full-screen NN textured quad is absolutely nothing for any GPU. And compared to the gunk that is required to render a full-screen of arbitrary text through any gui toolkit ever it's many orders of magnitude less effort.

Ideally this would only ever exist at full-resolution in the on-screen framebuffer memory which would also make it extremely cheap memory wise.

But at least initially I would be going through JavaFX so it will instead have to have multiple copies and so on. The reason to use JavaFX is for all the auxiliary but absolutely necessary fluff like clipboard and dnd operations. I don't really like tabbed terminals (I mean I want to use windows as windows to multitask, not as a stack to task switch) but that is one way to ameliorate the memory use multiplication this would otherwise create.

So to begin with it would be extremely fat but that's just an implementation detail and not a limitation of the design.

worth bothering?

Still mulling that over. It's still a lot of work even if conceptually it's almost trivial.

Sunday, 22 November 2015

OpenCL 2.1 + java = zcl 0.x?

I noticed Khronos released the OpenCL 2.1 spec recently so I spent this morning updating zcl to include all the functions.

Since I don't have a suitable implementation nothing is tested and there's probably some typos and so on. I found a few small bugs in the enum tables while I was there.

But what took most of the time was the property queries. Each OpenCL object type has one or more query functions but rather than implement them all I use a tagged query function which branches to the correct function entry point at the lowest level but shares all the rest of the code. But then I had to add some specialist variants, and specialisations for return types and overloaded parameters - it started to get unwieldy and a new query type on CLKernel meant it wasn't going to be enough anyway.

So I said fuck that for a joke and just redid the whole mechanism.

For the basic 5-parameter queries I still share most of the code but I now add any type-specific queries separately. To cope with the api and code bloat i distilled the java side interface down to only two entry points for each query:

    native <T> T getInfoAny(int type, int ctype, int param_name);
    native <T> T getInfoAnyV(int type, int ctype, int param_name);

The first is a scalar query and the second an array one. It just means it now has to box primitive return types for scalar queries which is unlikely to have any measurable performance impact but the Java helpers which wrap the above interfaces in type-friendly calls could always be replaced with native equivalents if it was an issue.

This let me merge some internal jni code and delete a lot of snot and I moved the re-usability to a different layer so that the more specific queries can share most of the code. For example this was the previous set of native interfaces on CLObject, and although this covered the kernel and program specific 6-argument queries like GetProgramBuildInfo() it was getting a bit messy.

    native long getInfoLong(int type, long subtarget, int param);
    native long[] getInfoLongA(int type, long subtarget, int param);
    native int getInfoInt(int type, long subtarget, int param);
    native byte[] getInfoByteA(int type, long subtarget, int param);
    native <T> T getInfoP(int type, long subtarget, int param, int ctype);
    native <T extends CLObject> T[] getInfoPA(int type, long subtarget, int param, int ctype);
    native long getInfoSizeT(int type, long subtarget, int param);
    native long[] getInfoSizeTA(int type, long subtarget, int param);

... It seemed like a good idea at the time.

The exposed interfaces remain the same (like getInfoString(param), getInfoInt(param), etc).

Given the complete lack of interest and because it needs some testing anyway I wont be releasing a zcl-0.6 just yet.

Friday, 6 November 2015

OpenCL 2.0 + Java = zcl 0.5, or ~= 1.0 beta

I spent a wet morning doing some clean up and packaging of another zcl build and just finished updating the home page and uploading the source.

Although I just bumped the revision, this is getting pretty close to a 1.0 release. It's still got a few missing bits but it's mostly because the documentation is a bit broken beyond the README. It is only compatible with Java 8.

The home page has more details but the big points are that it now garbage collects everything (with explicit override), the lambda interfaces (trivial though they are), dynamically links to libOpenCL, fills out the extension framework and implements some extensions, and supports cross-platform building of native code.

I had to add a small code-generator to make the dynamic linking practical but it relies on the strict formatting of cl.h and does nothing fancy

Now i've got cross platform sorted out i'll probably do all my work to this interface rather than jogamp/jocl because it's just nicer to use and easier to work with. This might not mean any more frequent updates but at least it should get tested more. But apart from not being able to get SVM working at all on my machine (sdk demo works, cut and pasted bits from demo, or any other thing i write - crash) i've encountered very few bugs anyway.

I've probably covered enough of the new stuff in the blog previously so probably wont have much to add, but the curious are welcome to ask.

Friday, 23 October 2015

Catching up

Not much going on so today here's a diary entry ...

I've added a few little things to ZCL - it's coming along quite nicely, I should probably work toward another release. I'm slowly adding the functional-like stuff into it, I decided to go with 'q.offer()' as the enqueue function, and 'ofXX()' as the factory methods. Together with the garbage collector support it does offer some interesting possibilities for code-reuse but i'm still experimenting with it in practice.

I needed to access a webcam so i added another api to jjmpeg (but i might move it) which just wraps v4l2 devices directly from the file descriptor (no library). Actually I did that a while ago but have slowly been filling it out as I needed more functionality. OTOH I started looking into the total snot-show that webcam access is on microsoft platforms and decided to give up - you can't even build the media-framework libraries with mingw-w64 as far as I can tell and you need vs of some form just to get the "system" headers. Ugh.

However I did find webcam-capture library which has already solved these problems. It's probably what i will look at as a fallback, but on linux efficiency is a bit low. The simplest webcam dump to JavaFX with my library generates almost no garbage (it provides static access directly to the driver buffers) and uses 50% less cpu time compared to using the low-level interfaces it provides despite a pretty expensive YUYV conversion step. The high level swing one is 4x the cpu overhead.

Along the way I found that was actually using videoInput "library" but I couldn't get that to link (cross compile at any rate, it should work with the ms sdk) - and in any event that just uses the directshow stuff which I had working ... i dunno, years ago. But the driver no longer works for the webcam i'm using and i'd have to buy it, ... so yeah that can wait.

And that ultimately led me to openimaj which probably would've saved me doing almost all that myself, although i would've needed to understand it anyway for OpenCL translation. And maven, ffs. But I guess I should at least have a look.

Discovery of useful software isn't that easy these days with so much noise - even if you go looking which I can't say I did ...

I've also been using a small amount of OpenGL and interoperating with OpenCL. It's become all a bit ... naff, and JOGL has been necessarily messed up to support all this nafficity. Pity I had to do this now rather than in a couple of months otherwise I would be looking at vulkan instead but with any luck that will be out soon enough to move to it before I need to get too far into GL (with the steammachines in november?). I'm sure microsoft will find some way to totally fuck-up it's cross-platform parts again though. My current thinking is that i will write a java binding for it once I have my hands on it (if only just-because) but I haven't looked into it all so far. But removing the static per-thread state should make it a lot saner.

For now I have some simple classes to do some off-screen rendering, and some OpenCL interop which is enough for what I want. Access to an output texture in JavaFX would be nice but it sounds like this is just not going to happen. Although one would expect a vulkan backend to be done at some point it will probably suffer the same hiding issues (well, with good reason I guess). If I really needed more performance i'd just use another toolkit - which is sad as that doesn't appear to be the intended vision of the javafx designers.

On that interop I had to fill out the extension mechanism in ZCL. I followed the prototype I'd created earlier. Currently each extension is provided by a different CLExtension class. It holds a pointer which in C-land is a function table resolved on the platform, and each platform object manages these. At first I was just going to use this as the mechanism for accessing extensions but it quickly becomes messy - you have to find the platform the object belongs to an in some cases this requires multiple queries (e.g. q.getDevice().getPlatform()). One approach I tried was to hide that by providing the extension methods directly on the object they extend - e.g. new CLContext or CLCommandQueue methods. These then manage looking up the extension and invoking the correct method for the given object. The details still need to be resolved but only once per object and it's all handled java-side.

There's a bit more behind the original mechanism than just code tidyiness for the extensions - they could potentially be loaded at runtime, or written separately from the core library. But on reflection how useless is that? The problem with this approach is each extension has it's own object - this is good and bad in that eventually you end up with a table required per context, queue, device, or whatever.

I think putting the extension methods on the target object is correct and after that the details don't really matter so much since it's an internal detail. But (on the fly design) I guess i should just maintain a CLPlatform reference on each object which can be extended and handle it that way. The extension objects will still be per-extension which keeps a cleaner namespace but they only need to be set per-platform which doesn't happen often. I'm pretty sure all objects have a 1:1 platform relationship, that would be the only thing to throw a spanner in the works; but the whole extension mechanism wouldn't work if that were the case.

Somewhere along this journey i came across some C++ code for something, I can't remember what it was. It was how I find most C++ code - it's been so over-engineered almost all of the actual lines of code is just boilerplate. The workings are so hidden I gave up trying to find it. It's just as bad in Java land where everyone wants to write a fucking framework before they even get started. Cut the bullshit and get to the point. C++ shows its heritage as being born from a time when "Software Engineering" was going to solve the worlds software problems by taking the programmer out of programming; using UML and CASE tools and auto-generating everything. It really shows, it is not a good language. This craziness was at it's peak just as I was going through uni and its done it's fair share of damage to the world and clearly continues to if abominations like C++ still exists.

That'll do for now.

Friday, 2 October 2015

OpenVX

Here's something I missed being out of the loop for the last good while: OpenVX.

Although it appears so did the rest of the industry? Another standard sent to rot?

I'll have to have a closer look though; if it is well thought out it will be worth borrowing a few ideas at the least as it's a problem i've come up with multiple solutions for myself and it's always useful to get a fresh perspective.

Or maybe it'll give me something to play with along with ZCL either as a binding and/or a pure Java (re)implementation for prototyping.

I wonder how far away vulkan is for the rest of us.

Update: Some movement has occurred on the HSA front recently. I suppose just some more announcements or partnerships or something. There's a lot of technology and software that needs to come together to get anywhere and it's taking it's time. I wonder if it'll ever get traction though. I guess it solves one of the problems with OpenCL - you actually have to program parallel software which is a skill few have the time to learn and fewer companies have the budget to fund. At least it has the potential to stay in the compiler layer, hidden from your typical code-monkey.

OpenCL notes

As tangentially related observation I've been working on some OpenCL stuff of late and I came across a problem I just couldn't get to run very fast on the dev GPU. I ended up resorting to using an OpenCL/CPU kernel and another queue but along the way i tried a couple of other interesting things.

One was using a native kernel written in Java - but this still requires a CPU driver / queue since the GPU driver i'm using doesn't support native kernels. ZCL has fairly simple interface to this:

public interface CLNativeKernel {

    /**
     * Native kernel entry point.
     *
     * @param args arguments, with any CLMemory objects replaced with
     * ByteBuffer.
     */
    public void invoke(Object[] args);
}

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

The JNI hides the details but it behaves the same way as the C code whereby any memory objects in the argument list are replaced by pointers (ByteBuffer here). Not sure if i'll keep the varargs prototype because it is inconsistent with every other enqueue function and only saves a little bit of typing. I'll review it when i look at the functional stuff i mentioned in the last post.

Which can be used efficiently and conveniently in Java 8:

  CLBuffer points = cl.createBuffer(0, 1024 * 4);
  q.enqueueNativeKernel((Object[] args)-> {
    // object[0] = ByteBuffer = points contents
    // object[1] = Integer = 12
    // object[1] = Integer = 24
  }, null, null,
  points, 12, 24);

Since my prototype harness didn't have a CPU queue until I subsequently added it my first approach was to emulate a native kernel using event callbacks and user events. It actually worked pretty well and was about the same runnning time, although it took a bit more (fairly straightforward) code to set up.

One approach I took was to have two queues - the primary 'gpu' queue where the work is done, and another one used for the memory transfer and rendezvous point.

  // setup CLEventLists and the user event

  // gpu part of work
  gpu.enqueueXXKernel(..., gpudone);
  // prepare for cpu component
  work.enqueueReadBuffer(..., gpudone, readdone);
  // prepare for gpu again
  work.enqueueWriteBuffer(..., cpudone, writedone);
  // make sure q is ready
  gpu.enqueueMarkerWithWaitList(writedone, null);

  memdone.setEventCallback(CL_COMPLETE, (CLEvent e, int status) -> {
    // do the work
    cpudoneevent.setUserEventStatus(CL_COMPLETE);
  });

In this case all the enqueue operations are performed at once and events are used to synchronise. This simplifies the callback code a little bit. Now i'm looking it it there's probably no need for the separate queue if the gpu queue is synchronised with it anyway. (like with most of these examples it is a summary of what i came up with, but not the full sequence of how i got there which explains some of the decisions).

This is a trivial approach to ensuring the 'gpu' queue behaves as the caller expects: that is, as if the work was performed in sequence on the queue and without having to pass explicit events. I'm using the read/write interfaces rather than map/unmap or otherwise mostly out of habit, but the data in question is quite small so it shouldn't make much difference either way.

And FWIW for this problem ... this approach or the java NativeKernel one actually runs a tiny bit quicker than using the OpenCL/CPU device let alone the GPU (all wall-clock time on the opencl q).

I had to make some small tweaks to the CLEventList code to make this all work and to tie properly into the garbage collection system. Mostly this was adding a CLEvent array rather than just using the pointer array and fixing the implementation. I kept the pointer array to simplify the jni lookup. I also had to have construction go through the same mechanism as the other CLObjects so they retain global reference uniqueness. This greatly simplifies (i.e. completely removes) reference tracking which is always nice with async callbacks. I think it should all "just work"; it does from the Java side - but i need to check from the OpenCL side of things whether actions like setUserEvent() adds an explicit reference bump.

This is a prime example of what HSA should solve well, but for now this is what i've got to work with.

I've been so busy with details i haven't had a chance to even look at any OpenCL 2.0, let alone OpenVX, HSA, or much else. And frankly the spirit is just not that willing of late. Spring is just the latest of a litany of excuses for that.

Thursday, 1 October 2015

OpenCL lambda enqueue

Just had a thought on an alternative api for CLCommandQueue in zcl. No this has nothing to do with lambda calculus in OpenCL.

An inconvenience in the current api is that all the enqueue functions take a lot of arguments, many of which are typically default values. This can be addressed using function overloading but this just adds additional inconvenience as there are also simply a lot of functions to overload. A related issue is things like extensions can add additional entry points which are object-orientedly resident on the queue object but placing them there doesn't necessarily fit.

And finally new compound operations need to be placed elsewhere but also fit a similar semantic model of enqueing a task to a specific queue.

So the thought is to instead to use java's lambda expressions to create queueable objects which know how to run themselves, and then at least the waiters/events parameter overload can be handled in one place.

So rather than:

// some compound task
  public void runop(CLCommandQueue q, CLImage src, CLImage dst,
      CLEventList waiters, CLEventList events) {
     ... enqueue one or more jobs ...
  }
  public void runop(CLCommandQueue q, CLImage src, CLImage dst) {
     runop(q, src, dst, null, null);
  }
  public void runop(CLCommandQueue q, CLImage src, CLImage dst,
      CLEventList event) {
     runop(q, src, dst, null, event);
  }

// usages
 runop(q, src, dst, waiters, events);
 runop(q, src, dst, events);
 runop(q, src, dst);

I can do:

// the interface
interface CLTask {
  public void enqueue(CLCommandQueue q, CLEventList w, CLEventList e);
}

// the creation (only one required)
 public CLTask of(CLImage src, CLImage dst) {
   return (q, w, e) -> {
     ... enqueue one or more jobs ...
   };
 }

// usages
 q.run(op.of(src, dst));
 q.run(op.of(src, dst), events);
 q.run(op.of(src, dst), waiters, events);

This could extend throughout the rest of the api so that for example a CLBuffer would provide it's own read task factories:

  public CLBuffer {

    public CLTask ofRead(byte[] target) {
      return (q, w, e) -> {
        q.enqueueReadBuffer(this, true, 0, target.length, target, 0, w, e);
      };
    }
  }

// usages
  q.run(buffer.ofRead(target));
  q.run(buffer.ofRead(target), events);
  q.run(buffer.ofRead(target), waiters, events);

vs

// typical usage (without overloading)
  q.enqueueReadBuffer(this, true, 0, target.length, target, 0, null, null);
  q.enqueueReadBuffer(this, true, 0, target.length, target, 0, null, events);
  q.enqueueReadBuffer(this, true, 0, target.length, target, 0, waiters, events);

I think this would provide a way to add the convenience of overloading without a method count explosion. But the real question is whether it would actually improve the api in any meaningful way or merely make it different. Probably at this point it's a tentative yes on that one for many of the same reasons lambdas are convenient such as encapsulation and reuse.

There are some issues of resolving state at point-of-execution and threads but these are already an issue with OpenCL code to some extent and definitely with lambdas in general.

One could keep going:

// the interface
interface CLTask {
  public void enqueue(CLCommandQueue q, CLEventList w, CLEventList e);

  public default void on(CLCommandQueue q) {
    enqueue(q, null, null);
  }
  public default void on(CLCommandQueue q, CLEventList w, CLEventList e) {
    enqueue(q, w, e);
  }
}

// usage
 buffer.ofRead(target).on(q);

Despite this having the benefit of layering in isolation above the base api I think it starts to get a little absurd and turns into "much of a muchness" deckchair shuffling.

Although this addition is probably useful:

// the interface
interface CLTask {
  public void enqueue(CLCommandQueue q, CLEventList w, CLEventList e);

  public default CLTask andThen(CLTask after) {
     return (q, w, e) -> {
        enqueue(q, w, e);
        after.enqueue(q, w, e);
     };
  }
}

// usage
 q.run( buffer1.ofRead(target).andThen(buffer2.ofWrite(target)) );

Actually I didn't really intend it as an outcome but this also becomes a lot more usable if the resources in questions are automatically reclaimable via gc as per my last post. Whole state and work spaces can be retained and reused through nothing more than a CLTask reference.

I think i've convinced myself of the utility now but either way it takes very little code to try it.

OpenCL garbage

I was working on some higher level containers for managing OpenCL stuff and came to the conclusion that I wanted to add automatic resource reclaimation to zcl - it was either that or fill a whole hierarchy of objects with reference counting. But reference counting is slow, error-prone, and a big mess to write in so it isn't at all attractive when there is an alternative.

I'd already done it in jjmpeg but i wasn't really keen on the way i implemented it there and wanted to see if i could come with a more streamlined solution. Like when I did it for jjmpeg I started with this article about JavaSE finalisation and using weak reference queues.

I think the solution I came up with will work ... and it turned out to be rather simple in the end.

Previously all CLObjects were a simple lightweight pointer handle with all the details passed to the C functions. They all have an init(pointer) constructor which was called directly from the JNI layer. Duplicate objects referencing the same resource were not an issue so I just let it happen. Well it's easy to break but if you treat objects like the C pointers they are and know that dangling references are possible then it's not unsolvable.

But for GC to work the references need to be unique. This is fairly easy to guarantee as the resources are just memory pointers - which are guaranteed to be unique and unchanging. So rather than the JNI layer invoke the constructors directly I just call a factory method with a type index which lets me move some of the code into Java - it isn't significantly simpler but it is more flexible.

For the reference queue to work properly I need to store them in a container anyway so this conveniently meshes with using a hashtable to uniqify the objects.

  static CLObject toObject(int ctype, long p) {
    CLObjectHandle h = referenceMap.get(p);

    if (h != null)
      return h.get();

    return classTable[ctype].newInstance(ctype, p);
  }

My first attempt passed the Class through (this is how i did it in JNI) but I changed it to an integer. It makes the JNI a bit easier and having the type as an integer simplifies the release call (OpenCL api isn't OO and has per-type release functions). Being able to identify the object fully using primitive types also lets me freely use them without polluting the reference tree; which is critically important when dealing with gc.

Now comes the bit which i fucked up in jjmpeg (well the biggest bit). Each object is represented by 4(!) classes. An autogenerated native abstract class which includes the static native method prototypes and a hand-written native concrete class which implements any type-specific dispose or construction semantics. Then there is an autogenerated abstract public class which includes all the autogenerated methods again - this time invoking all the methods on the native class after looking up the object pointer. And finally a hand-written public concrete class which includes constructors, helpers, and any other special cases where the details are better hidden.

This is just a lot of code - every public method on the "java" class ends up calling a native method on the "native" class so every method needs at least two implementations; . This was the main driver for ZCL simply using a single JNI implementation and foregoing this redundant juggling of the call stack just to insert the resource pointer into the call. In most cases in ZCL the public api is just the native method and it needs no redundant wrapper.

This time I just added a single general-purpose CLObjectHandle weak reference type which is used by all instances to track the native resource. It just holds the pointer (and the ctype) and implements the release. I just add one of these to each CLObject in one place.

  public abstract class CLNative {
    final long p;

    protected CLNative(long p) {
      this.p = p;
    }
...
  }

  public abstract class CLObject extends CLNative {
    final CLObjectHandle h;

    protected CLObject(int ctype, long p) {
      super(p);
      h = new CLObjectHandle(this, ctype, p);
    }

...
    static class CLObjectHandle extends WeakReference<CLObject> {
      long p;
      int ctype;

      CLObjectHandle(CLObject referent, int ctype, long p) {
        super(referent, referenceQueue);
        this.p = p;
        this.ctype = ctype;
        referenceMap.put(p, this);
      }

      void release() {
        if (p != 0) {
          map.remove(p);
          CObject.release(ctype, p);
          p = 0;
        }
      }
    }
...

  }

This and a bit of house-keeping is all that is required.

Having release be idempotent allows explicit release mechanisms to remain - for those cases where you can't afford to let the native resource management be at the whim of the garbage collector. For this reason i may also have to move the native pointer resolution in the JNI from a CLNative.p field lookup to resolving it via the handle. I need to investigate the cost of doing this first, and also whether explicit release like this will actually work in practice (e.g. if you release an object with more than one reference, does it fuck up?). Doing this would also let me use the correct integral type if I felt the need by just creating two different CLObjectHandle classes (32/64) and resolving sizes in the JNI code.

There is some potential problems where you resolve an object for the first time via a non-referencing api (for example clGetProgramInfo(CL_PROGRAM_CONTEXT) and the like) and then let the reference expire. But this shouldn't normally be a problem since you would have to get the context before creating the program and are going to be keeping it around for the lifetime of the program and thus only one xxRelease is every invoked. And this should normally hold for everything else too. If it turns out to be an issue I have mechanisms I can use to address it from adding an explicit object reference to the given objects (e.g. a CLContext to each CLProgram created), or adding phantom reference bumps on specific apis.

It's actually a devilishly difficult thing to test and verify: even once you know the exact reference counting semantics of every OpenCL api the interaction with the JVM will hide faults.

I haven't explored further but having unique objects and gc lets me freely cache local copies of resource handles for convenience or efficiency and so on. It really simplifies using the library enough as it is.

The next zcl release will include this as well as a couple of bug fixes and some other things which make it easier to use. Dunno when that might be though.

Thursday, 17 September 2015

empirically corrected-approximate integer division

Although i haven't been posting about it i've been continuing to poke around in bits and pieces of code. Well, a little bit.

I did a bit of OpenCL last week and that was pretty fun. I had enough time to really dig into optimising a particular routine and was down to inspecting the ISA output from the driver. Good stuff. The GCN isa is pretty foreign to me so I had to use small snippets to isolate operations of interest.

For example one construct that comes up repeatedly when parallelising code is using a divide and modulus operator when splitting up a non-work-sized job into work-group sized blocks.

  int block_size = info.block_size;
  for (int id=get_local_id(0); id < limit; id+=64) {
    int block_no = id / block_size;
    int block_index = id % block_size;
     // do work
  }

Where possible one just chooses a power of 2 so this is a simple shift and mask, or integer divide by a constant isn't too bad as it can usually be optimised by the compiler. But this problem required a dynamic block size that wasn't a power of 2.

The solution? Use floating point multiply the reciprocal which can be calculated efficiently or here off-line. The problem is that this introduces enough rounding error to be worthless without some more work.

I must admit I just found the solution empirically here: i had a limited range of possible values so I just exhaustively tested them all against a couple of guesses. Hey it works, i'm no scientist.

  float block_size_1 = info.block_size_1;
  for (int id=get_local_id(0); id < limit; id+=64) {
    int block_no = (int)(id * block_size_1 + 1.0f / 16384);
    int block_index = id - (block_no * block_size);
     // do work
  }

This replaces the many instruction integer division decomposition with a convert+mad+convert.

On some work-loads this was a 25% improvement to the entire routine and these 2 lines are in an inner loop of about 50 lines of code.

Well it's been fun to play at this level again - its ... mostly ... pointless going to this level but just adds to the toolkit and I enjoy poking. Maybe one day i'll have a job where it's useful.

I gave zcl a go on this as originally I was thinking of trying some OpenCL 2 stuff but I may not bother now. Given the lack of use/testing it was pretty much bug free but I started filling out the API with some more convenient entry points. I also decided to add some more java-array interfaces here and there: they're just too convenient and it hides the mess in the C even if they might not be the most efficient in all cases.

This is the sort of thing i'm talking about:

  float[] data = new float[] { 1, 2, 3, 4, 5 };
  CLBuffer buffer = cl.createBuffer(CL_MEM_COPY_HOST_PTR, data);
vs
  float[] data = new float[] { 1, 2, 3, 4, 5 };
  ByteBuffer bb = ByteBuffer.allocateDirect(data.length * 4).order(ByteOrder.nativeOrder());
  bb.asFloatbuffer().put(data);
  CLBuffer buffer = cl.createBuffer(CL_MEM_COPY_HOST_PTR, data.length * 4, bb);
It's only two fewer lines of code ... but yeah that shit gets old fast. The first is more efficient too because this is a native method and it avoids the copy. In the first case CL_MEM_USE_HOST_PTR throws an exception though, and in the second it works (library call permitting).

The main downside is adding these convenience calls blows out the method count very quickly if you support all the primitive types - which detracts from the ease of use they're supposed to increase.

Another release? Who knows when.

And this week i've been poking at some OpenGL. My it's grown. I'm experimenting using JOGL for this although i'm not a fan of some of it's binding choices. It's crossed my mind but i'm pretty sure i don't want to create yet another binding as in a 'ZGL'. Hmm, I wonder if vulkan will clean up the cross platform junk from opengl.

Unfortunately my home workstation seems to have developed a faulty DIMM or something (unrelated note of note).

Monday, 25 May 2015

sgemm, OpenCL

Yesterday I couldn't do much else so I played with some OpenCL code again. Just as my left foot was nearing recovery from gout ... I think I strained my right foot from too much walking or other activity and i'm immobile again. Argh.

With OpenCL still haven't managed to work out why I can't use SVM - I have a C test, a C test based on extracting all the relevant code from the BufferBandwidth sample (from amd sdk), and a C++ test based on the BufferBandwidth sample; they all crash as soon as I try to invoke a kernel against an SVM buffer, although BufferBandwidth runs fine. I even tried compiling linux 3.19.8 - I had to modify the catalyst driver a little bit to get it to build but I had it working for a bit, but suspend was broken and then I made one too many changes to the linux config and i couldn't get it to boot again and eventually just gave up. The linux config system is pretty shit and any changes force a full rebuild so i was getting sick of that. When I did have it running it made no difference to the SVM stuff anyway.

OTOH I did find that using CL_MEM_USE_HOST_PTR works in much the same way anyway (in terms of java usefulness) - even without mapping or unmapping the values are being updated on the via the GPU device, so with any luck map/unmap are just noops. I didn't really look too much further though.

What I looked at instead was implementing a basic matrix matrix multiply, i.e. lapack's sgemm. Not really for any need but just curiosity; how much effort is required vs the payoff.

My test case was a C=AxB where each is (1024,1024) with row-major order storage. CPU is a AMD A10-7850K (kaveri).

  java naive             20.5
  java copy col B         1.5
  java copy col B mt      0.50
  opencl cpu naive        6.5
  opencl cpu float4x4     0.49
  opencl gpu simple       0.26
  opencl gpu float4x4     0.045

  java ojalgo (double)    0.48
  java la4j (double)      1.7

  java copy stream        0.40
  java copy stream x4     0.30

java naive
This just implements the classic algorithm literally - i.e. for all rows of A, dot product of row by all columns of B, etc. The problem here is that each dot product scans a column of B in a potentially worst-case way in terms of cacheable memory access - this size hits that.
java copy col B
This just inverts the two outer loops so that it runs for all columns of B and then dot products that with all rows of A. It copies the current column of B in the outermost loop and so it only has to run once for every 2^20 accesses (in this case). Which is obviously worth it.
java copy col B mt
This replaces the outer loop with a IntStream.range(0, n).parallel().forEach(). It's not optimal memory-use wise but that makes little difference in this canned example (see the last couple of results). This is a trivial change and also easily worth it.
opencl * naive
This is a trivial opencl implementation that runs transposed with each work item calculating a single output value. The work size is 64,1,1 in each case. This shows that it isn't worth using OpenCL without a bit more effort on the algorithm.
opencl * float4x4
This is the most complex implementation whereby each work-item calculates a 4x4 output cell (calculates 4 rows at once). The number of columns and rows must be multiple of 4. It's basically just an unrolled loop using vector types; but the code is still quite straightforward. At least in this case - since the problem is embarrassingly parallel - the effort required is modest for the gains possible.

java copy stream
This replaces the outer loop of the copy col B loop with a custom non-gc-polluting spliterator over the columns of B. i.e. it allocates one row for each thread which is re-used for each call. This is moderately more work to set-up but the spliterator is re-usable. It's also possibly slightly misleading due to the way hotspot optimises callbacks.
java copy stream x4
Well what the hell - this unrolls the inner loop by 4x so only works on matrices with A_n a multiple of 4.

The opencl code is sub-optimal for the CPU case - something closer to the java implementation would make sense - I will try again at another time perhaps. I'm not that familiar with the compiler behavior or best processing model to use for the CPU driver but using vectors obviously helps. But if it's barely faster than Java there wont be much point. OTOH a 10x speedup using the kaveri GPU is a bit more interesting.

Sorry no code today - if i keep poking i might drop it into a zcl-samples thing later on. I'm sure there's plenty of (better) code out there in accelerated lapack libraries anyway.

Update: So before i posted this i came across the java matrix benchmark and the pretty simple 'java copy col b' is pretty close to ojAlgo running on this box. I only ran it on this same test and not the benchmark so i don't know if it's 'fast' on this machine although i imagine most of it's perf advantages in multiply in the benchmarks is from multi-threading. I also had some time to blow so I tried the row stream and row stream unrolled versions just out of curiosity.

ojalgo only seems to do double arrays, which doesn't make much difference with this problem size apart from double the storage space. I'm using a double accumulator for the dot product fwiw.

Update: Bored. Looked at la4j. It's maven only but a quick makefile fixed that abhorrence. Anyway it's just a tiny bit slower in this case and surprisingly only single-threaded (for something which is `current', this really is surprising). It's using the same algorithm as "java copy col B" but it uses 2D java arrays for it's dense matrix and creates a garbage copy of every column during operation rather than re-using the array. It looks like a pretty nice little library apart from a few odd looking decisions like these, especially the custom serialisation mechanism, and lack of threading. But there's a lot more to a matrix library than a multiply.

FWIW I didn't bother to include it in the above but on the weekend I also tried a direct ByteBuffer as storage. It takes a bit longer than the array backend for hotspot to optimise but it's quite close to the array version once it has. Or actually a bit faster in the mt case, for some strange reason.

Sunday, 10 May 2015

faster faster loops

Given i haven't touched opencl for a while I thought i'd stop faffing about with threads and streams and see what this APU can do.

But I found a silly bug in zcl which rendered it broken so just mucked around playing with that and got nowhere with my original aim ...

I call a bunch of JNIEnv *A() functions, these take an array of jvalue's which should presumably be more efficient than walking a varargs list (if insignificantly so). But in an effort to clean up the way I was using it I broke it and hadn't gotten around to actually running anything until now. I will drop another zcl at some point but considering nobody's downloaded it i'm in no rush.

I also worked out some issue with the GPU driver, and possibly slackware. The mandelbrot demo works fine with javafx but other non-gui code just wasn't finding the GPU device. I couldn't work out what was going on but a strange error indicated it was probably some xfce session setup thing. I found an acceptable workaround in just setting export COMPUTE=localhost:0.0.

And then i spent the rest of the evening trying to work out why SVM wouldn't work on the GPU. It "works fine" on the CPU driver but although other operations are fine any kernel invocation leads to an insta-crash. After re-checking every code path i came to the conclusion that it's not the way i'm calling it, it just doesn't want to work.

And just now I tried a stripped down C implementation and it just crashes when I invoke a (do-nothing) kernel with an SVM argument :-/ Blast. I checked the BufferBandwidth sample and once I figured out netbeans was sucking too much for it to run and closed it; it worked fine. After a pretty long look i can't see why it isn't working so i must've done something really silly and small.

One part of svm - the common address pointers - aren't as useful from Java as from C anyway but the ability to share buffers without explicit map/unmap calls in the fine grained case should be, particularly on this APU.

netbeans

Netbeans is really starting to struggle for some reason. I was doing a big cleanup of a prototype which the boss gave to the customer (sigh) and moving lots of code around and suddenly it decided I had no main class and wouldn't even compile. After cleaning caches and other junk it was 'just' a non-obvious parser error. But I still had to resort to emacs+makefile to go through the compile errors one by one until I could get it to run. And the line that got it working in netbeans again was just a reference to a deleted import - i'd moved it to a common library.

But then it just started scanning dozens of files (dozens of times each) every time i switch windows; pausing for 250-1000ms each time. Cleaning the cache made no difference and it's already on an SSD. And it's running out of memory constantly - which messes up the incremental compilation something fierce. The last thing I did was I tried the same thing I did at home and disabled a dozen or so plugins i'll never use; but it didn't make much long-term difference at home so i'm not confident it will at work either.

But opening the zcl projects back up at home has pretty much busted it here - it's constantly running out of memory and taking a second or so to save any file and often not compiling it. I mean, it's actually become unfit for purpose.

Maybe lambda parsing is throwing it for a loop; but why? Any parameter/type matching should be somewhat limited in scope unlike C++.

Wednesday, 18 March 2015

from link to dlsym call

I've been spending an inordinate amount of time the last few days playing with make and doing some work on zcl. Been staying up way too late and then doing it all again the next day. Been a bit of fun poking at little problems though.

The java makefile system i've been working on has got a little fiddly due to some features i'm trying to implement; but so far all problems have been solvable in a fairly straightforward way and it's mostly now taking time because i'm writing some fairly long article(s) about the work. Most of the complication is with generated source and resource files and dealing with the command line to jar which is like but not quite the same as tar.

So I took a break to play with zcl again. Because the sdk version of libOpenCL may not match the system version i've added an option to compile with dynamic linking - i think the linker can also do this automagically but it's something i'd need to learn about first. But just changing it over turned out to be pretty easy an infact most of the work is hidden in a single auto-generated file together with a few small bits of checking to throw exceptions when functions are not available.

I wrote a horrible bit of perl that parses the very strict format of cl.h and creates a list of static function pointers and an init function to set them up, as well as C macros which automagically change the source to go via the function pointers rather than via dynamic linking.

So for example it creates the following bits for a given function:

...
static cl_int (*pclBuildProgram)(cl_program, cl_uint, const cl_device_id *, const char *, void (CL_CALLBACK *)(cl_program, void *), void *); 
...
static void init_clfunctions(void) {
        ...
        pclBuildProgram = cl_symbol("clBuildProgram");
        ...
}
...
#define clBuildProgram (*pclBuildProgram)
...

And I just include this at the head of the file and call init during the JNI OnLoad. Saved most of the work since none of the rest of the code needs to change unless I want to add pointer null checks. This opens the possibility of a "sdk-free" build of zcl since without the need to link it can get everything it needs from the header files which I can bundle.

I also looked into a mechanism for implementing CL extensions - I think i've worked out a decent solution (~interfaces on the platform) but haven't (fully) implemented any yet. I'm having second thoughts on whether the array interfaces for asynchronous calls for buffers are really adding enough to counter the work required to implement them and the bloat of the api size, but for now they will remain.

It's not really so "simple binding for java" anymore at ~10KLOC, but then again most of the parts are simple on their own.

After I remembered how to get going in OpenCL I did a little bit of testing on some of the new interfaces in opencl 2.0 - so far they seem to function; although for some reason my catalyst driver is only reporting the CPU device so it's not the best test - i think i need to reboot the machine because of memory fragmentation. Netbeans has turned into (more of) a pig of late and a few hundred tabs in firefox tends to swallow (all) memory. I guess I made a mistake on the memory size in this machine.

Monday, 16 March 2015

zcl 0.4

Another thing over the weekend I did was update zcl to cover OpenCL 2.0 and drop out a 0.4 release. It's completely untested at this point.

See the homepage

Once I do some testing and perhaps revisit the libCL linking mechanism I guess I'll do a 1.0, but given 0.3 was a year ago, ... yeah - don't hold your breath :)

Wednesday, 10 September 2014

Boycott nvidia, cuda?

nvidia has taken the nuclear option to sue every other gpu maker in existence (apart from ati/amd with which they already have cross-licensing agreements i guess).

Patent trolling is usually the last gasp of a failing business. Which implies that despite (or because of) their overpriced hardware they are failing as a manufacturer; GPus are now commodity items and the margins no longer exist to run their type of high-margin business.

Patents are a cruel abomination which distort the workings of a "free market"; they directly codify rentier behaviour which costs society both economical and technological progress. The only beneficiaries are the unproductive leeches of society at the cost of everyone else.

If you're an engineer or scientist who is currently using or considering cuda for your work I suggest you reconsider both to protest this failure of a strategy and to protect the future value of your work.

Just for nvidia to consider this strategy shows they are not long for this world and choosing to use such a single-supplier would be foolhardy.

Friday, 21 March 2014

thoughts on opencl + array methods

I was going to have a quick look at removing the erroneous asynchronous Get/SetPrimitiveArrayCritical() stuff from zcl this morning but I've hit a complication too far for my tired brain.

I changed to just allocating a staging buffer and using Get/Set*ArrayRegion() for the read/writeBuffer commands. I only allocate enough memory for the transfer and copy the transfer size around and so on. It's a bit bulky but it's fairly straightforward.

Then I started looking at the image interfaces and realised doing the same thing is somewhat more complex - either I have to copy the whole array to/from the staging buffer to/from each time (if the get/set updates on a portion of the image for example) or I have to flatten the transfer myself. The former pretty much makes the function pointless and the latter bulks out the binding and may require lots of jvm calls.

So now i'm deciding whether I just force synchronous transfers for all array interfaces because they probably have some use despite synchronisation being the mind-killer, or just deleting them altogether to drop a ton of code. Since i've already got all the code the former will probably be the approach I take. The event callback stuff i'm using to finish up the transaction seems pretty expensive anyway so there may not be much net difference (against a net use count of zero for the library, at that - it's just something to pass the time).

On another note I decided to use cvs as my local repository backend to store this stuff. I think the all-day-never-finished checkout of gcc finally tipped me over but I never liked subversion because it's too slow and is just shit at merging. I was surprised netbeans detected it and offered to install the cvs plugin automatically (and i was a little surprised it was already installed too). I don't need or want to use tools that weren't designed for my use-case.

Ho hum, back to work Tuesday. My boss actually apologised for taking so long to get the contract sorted but yeah, i'm not complaining! Looks like it'll mostly be a continuation of one of the projects i'm not terribly keen on too. Bummer I guess. All I really care about right now is sleep though.

Update: (I kept poking) I just removed the async handling code and force a blocking call. Get/SetPrimitiveArrayCritical() is used to access the arrays directly. I'll do a release another day though.

Monday, 17 March 2014

sumatra, graal, etc.

I didn't really wanna get stuck building stuff all day but that seems to have happened.

Sumatra uses auto* so it built easily no problem. Bummer there's no javafx but hopefully that isn't far off.

graal is a bit of a pain because it uses a completely custom build/update/everything mega-tool written in a single 4KLOC piece of python. Ok when it works but a meaningless backtrace if it doesn't. Well it is early alpha software I guess. Still ... why?

Anyway ... I tried building against 'make install' from sumatra but that doesn't work, you need to point your JAVA_HOME at the Sumatra tree as the docs tell you to. My mistake there.

So it turns out getting the hsail tools to build had some point after-all. The hsailasm downloaded by the "build" tool (in lib/okra-1.8-with-sim.jar) wont work against the libelf 0.x included in slackware (might be easier just building libelf 1.x). So I added the path to the hsailasm I built myself ... and ...

 export PATH=/home/notzed/hsa/HSAIL-Instruction-Set-Simulator/build/HSAIL-Tools:$PATH
 ./mx.sh  --vm server unittest -XX:+TraceGPUInteraction \
   -XX:+GPUOffload -G:Log=CodeGen hsail.test.IntAddTest
[HSAIL] library is libokra_x86_64.so
[HSAIL] using _OKRA_SIM_LIB_PATH_=/tmp/okraresource.dir_7081062365578722856/libokra_x86_64.so
[GPU] registered initialization of Okra (total initialized: 1)
JUnit version 4.8
.[thread:1] scope: 
  [thread:1] scope: GraalCompiler
    [thread:1] scope: GraalCompiler.CodeGen
    Nothing to do here
    Nothing to do here
    Nothing to do here
    version 0:95: $full : $large;
// static method HotSpotMethod
kernel &run (
        align 8 kernarg_u64 %_arg0,
        align 8 kernarg_u64 %_arg1,
        align 8 kernarg_u64 %_arg2
        ) {
        ld_kernarg_u64  $d0, [%_arg0];
        ld_kernarg_u64  $d1, [%_arg1];
        ld_kernarg_u64  $d2, [%_arg2];
        workitemabsid_u32 $s0, 0;
                                           
@L0:
        cmp_eq_b1_u64 $c0, $d0, 0; // null test 
        cbr $c0, @L1;
@L2:
        ld_global_s32 $s1, [$d0 + 12];
        cmp_ge_b1_u32 $c0, $s0, $s1;
        cbr $c0, @L12;
@L3:
        cmp_eq_b1_u64 $c0, $d2, 0; // null test 
        cbr $c0, @L4;
@L5:
        ld_global_s32 $s1, [$d2 + 12];
        cmp_ge_b1_u32 $c0, $s0, $s1;
        cbr $c0, @L11;
@L6:
        cmp_eq_b1_u64 $c0, $d1, 0; // null test 
        cbr $c0, @L7;
@L8:
        ld_global_s32 $s1, [$d1 + 12];
        cmp_ge_b1_u32 $c0, $s0, $s1;
        cbr $c0, @L10;
@L9:
        cvt_s64_s32 $d3, $s0;
        mul_s64 $d3, $d3, 4;
        add_u64 $d1, $d1, $d3;
        ld_global_s32 $s1, [$d1 + 16];
        cvt_s64_s32 $d1, $s0;
        mul_s64 $d1, $d1, 4;
        add_u64 $d2, $d2, $d1;
        ld_global_s32 $s2, [$d2 + 16];
        add_s32 $s2, $s2, $s1;
        cvt_s64_s32 $d1, $s0;
        mul_s64 $d1, $d1, 4;
        add_u64 $d0, $d0, $d1;
        st_global_s32 $s2, [$d0 + 16];
        ret;
@L1:
        mov_b32 $s0, -7691;
@L13:
        ret;
@L4:
        mov_b32 $s0, -6411;
        brn @L13;
@L10:
        mov_b32 $s0, -5403;
        brn @L13;
@L7:
        mov_b32 $s0, -4875;
        brn @L13;
@L12:
        mov_b32 $s0, -8219;
        brn @L13;
@L11:
        mov_b32 $s0, -6939;
        brn @L13;
};

[HSAIL] heap=0x00007f47a8017a40
[HSAIL] base=0x95400000, capacity=108527616
External method:com.oracle.graal.compiler.hsail.test.IntAddTest.run([I[I[II)V
installCode0: ExternalCompilationResult
[HSAIL] sig:([I[I[II)V  args length=3, _parameter_count=4
[HSAIL] static method
[HSAIL] HSAILKernelArguments::do_array, _index=0, 0xdd563828, is a [I
[HSAIL] HSAILKernelArguments::do_array, _index=1, 0xdd581718, is a [I
[HSAIL] HSAILKernelArguments::do_array, _index=2, 0xdd581778, is a [I
[HSAIL] HSAILKernelArguments::not pushing trailing int

Time: 0.213

OK (1 test)

Yay? I think?

Maybe not ... it seems that it's only using the simulator. I tried using LD_LIBRARY_PATH and -Djava.library.path to redirect to the libokra from the Okra-Interface-to-HSA-Device library but that just hangs after the "base=0x95..." line after dumping the hsail. strace isn't showing anything obvious so i'm not sure what's going on. Might've hit some ubuntu compatibility issue at last or just a mismatch in versions of libokra.

On the other hand ... it was noticeable that something was happening with the gpu as the mouse started to judder, yet a simple ctrl-c killed it cleanly. Just that alone once it makes it into OpenCL will be worth it's weight in cocky shit rather than just hard locking X as is does with catalyst.

Having just typed that ... one test too many and it decides to crash into an unkillable process and do weird stuff (and not long after I had to reboot the system). But at least that is to be expected for alpha software and i've been pretty surprised by the overall system stability all things considered.

I think next time I'll just have a closer look at aparapi because at least I have that working with the APU and i'm a bit sick of compiling other peoples code and their strange build systems. Sumatra and graal are very large and complex projects and a bit more involved than I'm really interested in right now. I haven't used aparapi before anyway so I should have a look.

If the slackware vs ubuntu thing becomes too much of a hassle I might just go and buy another hdd and dual-boot; I already have to multi-boot to switch between opencl+accelerated javafx vs apu.

Update: Actually there may be something more to it. I just tried creating my own aparapi thing and it crashed in the same way so maybe i was missing some env variable or it was due to a suspend/resume cycle.

So I just had another go at getting the graal test running on hsail and I think it worked:

 export JAVA_HOME=/home/notzed/hsa/sumatra-dev/build/linux-x86_64-normal-server-release/images/j2sdk-image
 export PATH=${JAVA_HOME}/bin:${PATH}
 export LD_LIBRARY_PATH=/home/notzed/hsa/Okra-Interface-to-HSA-Device/okra/dist/bin
 ./mx.sh  --vm server unittest  -XX:+TraceGPUInteraction -XX:+GPUOffload -G:Log=CodeGen hsail.test.IntAddTest

[HSAIL] library is libokra_x86_64.so
[GPU] registered initialization of Okra (total initialized: 1)
JUnit version 4.8

...

Time: 0.198

OK (1 test)

Still, i'm not sure what to do with it yet ...

I was going to have a play today but I just got the word on work starting again (I can probably push it out to Monday) so I might just go to the pub or just for a ride - way too nice to be inside getting monitor burn. I foolishly decided to walk into the city yesterday for lunch with a mate I haven't seen for years (and did a few pubs on the way home - i was in no rush!) but just ended up with a nice big blister and sore feet for my troubles. It's about a 45 minute walk into the city but I don't do much walking.

I want see if anything interesting comes out of the Sony's GDC talk first though (the vr one? - yes the VR one, it's just started).

And ... done. Interesting, but still early days. Even if they release a model for the public at a mass-market price it's still going to have to be a long term project. Likely the first 5-10 years will just be experimentation and getting the technology the point where it is good and cheap enough.

Update: Yep, i'm pretty sure it's just a problem with suspend/resume. I just tried running it after a resume and it panicked the kernel.

hsa, gcc hsail + brig

I didn't spend much time on it yesterday but most of it was just reading up on hsa. There's only a few slides plus a so-shittily-typeset-is-has-to-be-microsoft-word reference manual on hsail and brig. As a bonus it includes some rather poor typeface choices (too stout - the trifecta of short, thick, and fat), colours, and tables to complement the word-wrap that all say to me 'not for printing' (not that I was going to). The link to the amd IOMMUv2 spec is broken - not that i'm likely to need that. Bit of a bummer as amd docs are usually formatted pretty well and they seem to be the primary driver at this point. In general the documentation 'needs work' (seems to be my catchphrase of 2014 so far) although new docs and tools seem to be appearing in drips and drops over time.

I'd already skimmed some of it and had a general understanding but I picked up a few more details.

The queuing mechanism looks pretty nice - very simple yet able to do everything one needs in a multi-core system. I had been under the impression that the queue system had / needed some hardware support on the CPU too but looking at it but it doesn't look necessary so was just a misunderstanding. Or ... maybe it's not since any work signalling mechanism would ideally avoid kernel interactions and/or busy waiting - either or both of which would be required for a purely software implementation. But maybe it's simpler than that

To be honest i would have preferred that hsail was a proper assembly language rather than wrapping the meta-data in a pseudo-C++ syntax. And brig seems a bit unnecessarily on the bulky side for what is essentially a machine code encoding. At the end of the day neither are deal breakers.

The language itself is kind of interesting. Again I thought it was a slightly higher level virtual-processor that it is, something like llvm's intermediate representation or PTX. But it has a fixed maximum number of registers and the register assignment and optimisation occurs at the compiler stage and not in the finaliser - looks a lot more like say DEX than IR or Java bytecode. This makes a lot of sense unless you have a wildly different programming model. Seems a pretty reasonable and pragmatic approach to a universal machine code for modern processors.

The programming and queuing model looks like something that should fit into Epiphany reasonably well. And something that can be used to implement OpenCL with little work (beyond the compiler, but there's few of those already).

GCC

I managed to get gcc checked out to build. The hsa tools page just points to the subversion branch with no context at all ... but after literally 8 hours trying to check it out and only being part-way through the fucking C++ standard library test suite, I gave up (I detest git but I'm no fan of subversion by any stretch of the imagination). I had to resort to the git mirror. Unfortunately gcc takes a lot longer to build than last time I had to despite having faster hardware, but that's 'progress' for you (no it's not).

I'm not sure how useful it is to me as it just generates brig directly (actually a mash up of elf with amd64 + brig) and there's no binutils to play with hsail that I can tell. But i'll document the steps I used here.

  git clone --depth 1 -b hsa git://gcc.gnu.org/git/gcc.git

  mkdir build
  cd build
  ../gcc/configure --disable-bootstrap --enable-languages=c,c++ --disable-multilib
  make

Slackware 64 is only 64-bit so I had to disable multilib support.

The example from gcc/README.hsa can then be compiled using:

  cd ..
  mkdir demo
  cd demo
  cat > hsakernel.c
extern void square (int *ip, int *rp) __attribute__((hsa, noinline));
void __attribute__((hsa, noinline)) square (int *in, int *out)
{
  int i = *in;
  *out = i * i;
}
CTRL-D
  ../build/gcc/xgcc -m32 -B../build/gcc -c hsakernel.c  -save-temps -fdump-tree-hsagen

Using -fdump-tree-hsagen outputs a dump of the raw HSAIL instructions generated.

[...]

------- After register allocation: -------

HSAIL IL for square
BB 0:
  ld_kernarg_u32 $s0, [%ip]
  ld_kernarg_u32 $s1, [%rp]
  Fall-through to BB 1

BB 1:
  ld_s32 $s2, [$s0]
  mul_s32 $s3, $s2, $s2
  st_s32 $s3, [$s1]
  ret_none 

[...]

Went through the gcc source and found a couple of useful bits. To get the global work-id I found you can use: __builtin_omp_get_thread_num() which compiles into workitemabsid_u32 ret,0. And __builtin_omp_get_num_threads() which compiles into gridsize_u32 ret,0. Both only work on dimension 0. And that seems to be about it for work-group functions.

I'm not really sure how useful it is and unless the git mirror is out of sync there hasn't been a commit for a few months so it's hard to know it's future - but it's there anyway.

My understanding is that a reference implementation of a finaliser will be released at some point which will make BRIG a bit more interesting (writing one myself, e.g. for epiphany, is a bigger task than i'm interested in right now). I'm probably going to have more of a look at aparapi and the other java stuff for the time being but eventually get the llvm based tools built as well. But ugh ... CMake.