nengo / nengo-ocl Goto Github PK

Something about the recent learning changes broke PES in Nengo OCL. Here's what I get running learn_communication_channel (which works in Nengo reference):

I'll investigate further.

Nengo has indexable signals now. Support them.

Also, nengo has the new registry for neuron types and stuff. Support that too.

Hello,

I wanted to use Nengo Ocl in a virtual envrionnement but importing the module failed.
Here is the error message I got in python shell:

import nengo_ocl
:0: FutureWarning: IPython widgets are experimental and may change in the future.
Traceback (most recent call last):
File "", line 1, in
File "/home/menelou/MasterArbeit/Code/Nengo/nengo_ocl/nengo_ocl/nengo_ocl/init.py", line 7, in
from .sim_ocl import Simulator
File "/home/menelou/MasterArbeit/Code/Nengo/nengo_ocl/nengo_ocl/nengo_ocl/sim_ocl.py", line 11, in
from nengo.processes import StochasticProcess
ImportError: cannot import name StochasticProcess

Indeed in the last commit of Nengo, there is no class nor function in processes.py.
Or did I do something wrong ?
Thanks for your help :)
Menelik

In trying to use OCL on a mac, I got this.

RuntimeError: clBuildProgram failed: build program failure -

Build on <pyopencl.Device 'Intel(R) Core(TM) i7-3720QM CPU @
2.60GHz' on 'Apple' at 0xffffffff>:

No kernels or only kernel prototypes found when build executable.
(options: -I
/Users/celiasmi/.virtualenvs/__ip/lib/python2.7/site-__packages/pyopencl/cl)
(source saved as
/var/folders/91/__dqhtq4ts7vd187fs23ctsg340000gn__/T/tmpLJbQLJ.cl)


ERROR: An unexpected error occurred while tokenizing input
The following traceback may be corrupted or invalid
The error message is: ('EOF in multi-line string', (1, 8))


Trevor suggested I run the test_sim_ocl script, which i did, and it
failed (though it detected my cpu)...

So then I went here:
http://wiki.tiker.net/__PyOpenCL/Installation/Mac

<http://wiki.tiker.net/PyOpenCL/Installation/Mac>

And followed the steps at the bottom (for mac os x lion), and reran
the tests and got the same thing. 

OS is Lion, 10.9

Semantically it makes sense to filter signals. Giving each signal its own pstc provides necessary flexibility.

This addresses a point raised by @hunse in #17 .

the clra_gemv planner analyzes the geometry of the dot products and splits off special cases. All of these special case kernel launches can be run at once.

Why is misc_gemv1 slower than the lif kernel on the 7970 even when the inner product (the for loop) is removed ?? Just writing 0 to Y_data seems to take longer than the lif kernel. (Looking at test_one_random)

If the caller provides a Python float for these arguments, that should mean it is a broadcasted constant.

It is tempting to interpret numpy arrays as constants, since they cannot be directly modified by OpenCL kernel calls, and opens the door to analyzing patterns in the constants for customized code generation.

Some models call for neurons to have some noise injected to their membrane potentials on every timestep.

The simplest thing to do is to call a special kernel for generating the random numbers, and add them to the bias on every timestep. Consider adding the random number generation into the LIF kernel itself.

PyOpenCL provides a RANLUX generator.

Look at how random numbers are generated by Theano's MRG generator, how they are generated by reika, and how they are generated by PyOpenCL's RANLUX implementation.

@hunse Told me to open a discussion on this repo to discuss potential improvements.

Firstly, I'd like to say that Multi GPU may become extremely viable with NVIDIA Pascal's release of NVLink, and AMD Polaris's Coherent Interconnect Fabric (I believe that's what the name is) as it may be much more practical that it was on previous generations to split the simulation of different ensembles between GPUs. I developed a similar system about 18 months ago for doing multi-agent clustering, and essentially performed a "Pseudo Boosting" algorithm on the CPU in order to determine which clusters to group together on the GPUs; however, due to the topology of a Nengo networking being known beforehand it is most likely possible to implement a similar system with significantly less overhead.

Secondly, I noticed that it appears OCL doesn't properly reuse memory that it already allocated on previous kernel executions and doesn't avoid branching

While this is fine currently, @hunse had mentioned that he perhaps wants to implement dynamic parallelism. When dynamic parallelism is in use, you need to have a much tighter control on memory and branching in order to avoid major bottlenecking

An example of well implemented dynamic parallelism is available here. Note that this file has a few syntax errors, but the idea is there. Roughly, no branch has an uneven number of instructions, I do not create new stack frames, and I try to reuse as much memory as possible. Furthermore, I do not recommend trampolining of any sort when doing dynamic parallelism. It isn't designed for that and you will run into limitations very very quickly.

Last time I used OpenCL, you couldn't assign more than one thread to an individual GPU. It is possible that Vulkan has changed this. Said being, BEFORE attempting to tackle dynamic parallelism, we should look into optimizing our method of batching tasks to the GPU from the CPU. Since the order of instructions is known at all points of the simulation, if we rework how we're managing memory a bit its very possible that we could well exceed the performance of dynamic parallelism, although there will probably be some performance boost over its existing state. If we see that we still need dynamic parallelism though, we can always try that too.

There are some issues that need to be brought up. DP is only supported on NVIDIA cards that support compute model 3.5 and up. Namely, 750Ti and >= GTX 770. DP is supported on all AMD GPUs 6xxx and up I believe (With a few odd and random exceptions though). Which brings up the concern that we can't move all of Nengo to DP, since GPUs of that caliber may not be common among members of the community. Also, I don't believe DP can execute kernels on other GPUs. I believe it is only limited to running kernels internally. Note trampolining is possible between GPUs. Don't do this. Don't. Srsly. You will regret it.

If we were to add Multi GPU support with limited CPU communication, I believe OpenCL has a way to do P2P GPU-to-GPU communication without CPU intervention. This would probably be the way we'd want to go. Cuda Equivalent. If we use a UVA system, then we don't need to worry about this; however, UVA is only useful if you don't have an in depth understanding of how the GPUs are communicating. Luckily nengo has a very well defined set of rules for this, pertaining to the topology of the network :P

There was a feature with DirectCompute back in DX11.1 that I only recently saw exists in Vulkan. We can premake the list of operations that the GPU needs to execute, as well as specify memory transfers between peers, and to the host. We can then execute this premade list long after the fact. (Eg: Even between sessions). This may also be useful since it has some of the same performance benefits of DP and multithreading without the need for either (Note: It technically does both slightly worse, but it has much much lower requirements for the end user as well as less CPU overhead.)

Finally, something I noticed when going through OCL. OpenCL has APU support, and Skylake has Iris Pro APUs and AMD's Zen will have perhaps equally as strong APUs. None of OCL is optimized for APUs. Perhaps we should include a mode for this? APUs have the benefit that Host to Device and Device to Host transfers are practically free. I imagine an APU mode would try to offload as much as it could to AVX/SSE while leaving the rest to the APU itself. This would be particularly useful for laptops and/or ARM devices.

Thanks
Louis

https://github.com/ctn-waterloo/nengo_theano/blob/master/nengo/scripts/abs_val.py

So I'm trying to get learning to run on the GPU with nengo_ocl. Basically I've taken what @tbekolay did in nengo/nengo#202 and made it do the same thing using signals and ops, so that we won't have to make any changes on the nengo_ocl side. What I have right now in nengo/builder.py is:

    @builds(nengo.nonlinearities.PES)
    def build_pes(self, pes):
        activities = pes.connection.pre.neurons.output_signal
        error = pes.error_connection.output_signal
        scaled_error = Signal(np.zeros(error.shape))
        shaped_scaled_error = SignalView(scaled_error, (error.size, 1), (1,1), 0)
        shaped_activities = SignalView(activities, (1, activities.size), (1,1), 0)

        decoders = pes.connection.decoder_signal
        lr_signal = Signal(pes.learning_rate * 0.001)

        self.model.operators.append(
                Reset(scaled_error))
        self.model.operators.append(
                DotInc(lr_signal, error, scaled_error, tag="PES: scale error"))
        self.model.operators.append(
                ProdUpdate(shaped_scaled_error, shaped_activities, 
                           Signal(1), decoders, tag="PES: update"))

I'm using SignalViews there because the learning algorithm requires an outer product, and the error and activities signals are both 1-D arrays (their shape attr has length 1). This gives the same results as nengo/nengo#202 with both the reference implementation and the nengo_ocl numpy simulator, but doesn't work with the ocl simulator. So two questions:

1. Does using a SignalView like this make sense, is that how they were meant to be used?
2. Anyone have any idea why it isn't working with the OCL simulator? Not really up to speed on how that thing works. Learning with OCL looks like this, if it helps at all:
   

https://github.com/ctn-waterloo/nengo_theano/blob/master/nengo/scripts/TRN.py

There are some things (e.g. data read-out, sync between multiple platforms) that don't have to happen as fast as synapses operate (~1ms) but still regularly and somewhat fast (~10ms). Consider formalizing this mechanism so that the "make_probe" functionality of Network can be ported to the OpenCL code.

Related: video input at e.g. 25 or 50 fps changes video buffer far less frequently than the synapse rate.

many_dots kernel should only load geometry for the matrices currently in use. Lower shared memory -> better occupancy.

Operators in Nengo have simplified significantly (most importantly, there is no longer a ProdUpdate). This means that MultiProdUpdate is mostly doing DotIncs, with some Copies in there.

One possibility would be to do copies separately, and replace MultiProdUpdate with something specific to DotIncs.

One advantage of MultiProdUpdate is it has some code to compress operations. For example, multiple dots writing to the same destination could be done together (the GEMV kernels have code for this), copies could be done along with dots to the same destination, etc. It is unclear if any of this compression is happening right now, since the graph sorting will try to make sure operations writing to the same place are not scheduled concurrently, I think.

Things to be simplified:

• remove as_update parameter, since it is unused.

Hello
for a project I have to use nengo_ocl. I tried to use the example code from the readme but it does not work. The line 'm = nengo.Model('foo')' fails (AttributeError: 'module' object has no attribute 'Model'), so I used the example 'Addition'. I replaced the simulator in the following way:

import nengo_ocl
import pyopencl as cl
ctx = cl.create_some_context()
sim = nengo_ocl.Simulator(model, context=ctx)

This is the result when running the python file:
File "NengoProject.py", line 52, in
sim = nengo_ocl.sim_ocl.Simulator(model, context=ctx)
File "/home/nengo_ocl-master/nengo_ocl/sim_ocl.py", line 53, in init
self, model=model, dt=dt, seed=seed, builder=builder)
File "/home/nengo_ocl-master/nengo_ocl/sim_npy.py", line 538, in init
builder = nb.Builder(copy=True)
TypeError: object() takes no parameters

Can you help me please?

It's a bit of a strange edge case, but I've often found it useful to have a Node with no inputs and no outputs to do some sort of timing operation. In nengo_gui this is used to keep javascript up to date with what the current simulator time is.

This sort of Node seems to crash nengo_ocl:

import nengo

model = nengo.Network()
with model:
    def timer(t):
        pass
    nengo.Node(timer, size_in=0, size_out=0)

import nengo_ocl
sim = nengo.Simulator(model)
sim.run(1)

When running the Nengo script 1-communication.py from @tcstewar 's repository nengo_benchmarks with the nengo_ocl backend, the following nengo_ocl kernel functions do not vectorize,

• probes
• gemv_many_dots
• elementwise_inc
• cl_lif_2
• filter_synapse

The spec for the machine being run on as well as the environment are as follows:

• 64 64-bit Intel Xeon CPUs
• CentOS 6.6
• PyOpenCL 2015.1
• OpenCL 1.2 driver
• Python 2.7
• Nengo 2.0

After setting PYOPENCL_COMPILER_OUTPUT, I collected output from the terminal and will email it to @hunse .

This is an obvious and easy place to ensure aligned reads. Round the number of structure variables up to the nearest (16? - what OCL device parameter is this?) before allocating.

Maybe my system is messed up, but running py.test on the current nengo_ocl master branch with Nengo 2.0.1 gives me a whole bunch of these attribute errors:

_________________________________________________ test.nengo.spa.tests.test_cortical.test_translate __________________________________________________

Simulator = <class 'nengo_ocl.sim_npy.Simulator'>, seed = 779372806

    def test_translate(Simulator, seed):
        with spa.SPA(seed=seed) as model:
            model.buffer1 = spa.Buffer(dimensions=16)
            model.buffer2 = spa.Buffer(dimensions=32)
            model.input = spa.Input(buffer1='A')
            model.cortical = spa.Cortical(spa.Actions('buffer2=buffer1'))

        output, vocab = model.get_module_output('buffer2')

        with model:
            p = nengo.Probe(output, 'output', synapse=0.03)

>       sim = Simulator(model)

../../../.pyenv/versions/2.7.9/lib/python2.7/site-packages/nengo-2.0.1-py2.7.egg/nengo/spa/tests/test_cortical.py:64: 
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ 
nengo_ocl/sim_npy.py:489: in __init__
    self._plan.extend(self.plan_op_group(op_type, op_list))
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ 

self = <nengo_ocl.sim_npy.Simulator object at 0x7f6c7bafcd50>, op_type = <class 'nengo.builder.synapses.SimSynapse'>
ops = [<nengo.builder.synapses.SimSynapse object at 0x7f6c7bb10d10>]

    def plan_op_group(self, op_type, ops):
>       return getattr(self, 'plan_' + op_type.__name__)(ops)
E       AttributeError: 'Simulator' object has no attribute 'plan_SimSynapse'

nengo_ocl/sim_npy.py:503: AttributeError

I just tried to run this model:

import nengo
import numpy as np

model = nengo.Network()
with model:
    stim = nengo.Node(lambda t: np.sin(t*2))

    ens = nengo.Ensemble(100, 1)
    nengo.Connection(stim, ens)

    def output(t, x):
        print t, x
    out = nengo.Node(output, size_in=1)
    nengo.Connection(ens, out)

import nengo_ocl
sim = nengo_ocl.Simulator(model)
sim.run(1)

and I get this error:

Traceback (most recent call last):
  File ".\test.py", line 17, in <module>
    sim = nengo_ocl.Simulator(model)
  File "c:\users\terry\documents\github\nengo_ocl\nengo_ocl\sim_ocl.py", line 57, in __init__
    self, network=network, dt=dt, seed=seed, model=model)
  File "c:\users\terry\documents\github\nengo_ocl\nengo_ocl\sim_npy.py", line 456, in __init__
    op_groups = planner(operators)
  File "c:\users\terry\documents\github\nengo_ocl\nengo_ocl\sim_npy.py", line 282, in greedy_planner
    assert chosen
AssertionError

Am I doing something wrong? It's been a while since I used nengo_ocl. Also, which version of nengo should I be using? I went back to 8dbe396 since master doesn't work with nengo_ocl at the moment....

Building a model in nengo_ocl takes about 10-15x longer than with regular nengo.
Here are some build times for differently sized models:

283950 neurons - nengo_ocl: 500s, nengo: 50s
488850 neurons - nengo_ocl: 1150s, nengo: 70s
892050 neurons - nengo_ocl: 3000s, nengo: 120s

Times were collected using a windows 8 machine with 2 x GTX660TI graphics cards, and an Intel i7-3770 3.40GHz quad-core processor.

Make sure it is correct!

Hello,

Nengo_ocl is working fine with standard synapses but my project aims to use also band pass filters synapses.
Taking the Github version of Nengo_ocl (the one from pip does not handle that of filter for synapses), I got one error from the plan_SimSynapse function which (if I good understood the code of that function) lies on make_step function of Nengo.
The later is not implemented in Nengo and I am wondering if there is a way to implement it knowing I have only bandpass filters.
(Maybe like the LinearFilter class from synapse.py ?)

Thanks in advance for your tips,

Ménélik

Currently the first step function in sim_ocl.Simulator takes a large fraction of a second, due perhaps to memory allocation on demand by OCL? Figure out what's up with that and ideally do it pre-emptively at simulator-construction time.

Think about how to handle custom_transforms... current way doesn't really count.

Some tests are failing on my CPU, compiling with the AMD APP. None of the Nengo tests, just one in test_ast_conversion.py, one in test_clra_gemv.py, and two in test_clra_nonlinearities.py.

Hello everybody,

As I was told in the previous thread, nengo_ocl does not support AdaptiveLIFRate.
I looked the code on the github and I came up with this modifications to support this neuron class (based on the code for LIFRate in nengo_ocl and AdaptiveLIFRate).

Could someone correct me if I am wrong (I hope not ^^) ?

(nb: the modifications did not return any errors as I tested my own model)

Here are all the modifications I made :

• In sim_ocl.py

from nengo_ocl.clra_nonlinearities plan_adaptive_lif_rate
from nengo.neurons import AdaptiveLIFRate```

def _plan_AdaptiveLIFRate(self, ops): 
    J = self.all_data[[self.sidx[op.J] for op in ops]]
    R = self.all_data[[self.sidx[op.output] for op in ops]]
    ref = self.RaggedArray(
        [np.array(op.neurons.tau_ref, dtype=J.dtype) for op in ops])
    tau = self.RaggedArray(
        [np.array(op.neurons.tau_rc, dtype=J.dtype) for op in ops])
    dt = self.model.dt
    tau_adapt = self.RaggedArray(
        [np.array(op.neurons.tau_n, dtype=J.dtype) for op in ops])
    inc = self.RaggedArray(
        [np.array(op.neurons.inc_n, dtype=J.dtype) for op in ops])

    return [plan_adaptive_lif_rate(self.queue, J, R, ref, tau, dt, tau_adapt, inc,
                              n_elements=2)]

def plan_SimNeurons(self,all_ops): 
    groups = groupby(all_ops, lambda op: op.neurons.__class__)
    plans = []
    for neuron_class, ops in groups:
        if neuron_class is LIF:
            plans.extend(self._plan_LIF(ops))
        elif neuron_class is LIFRate:
            plans.extend(self._plan_LIFRate(ops))
        elif neuron_class is AdaptiveLIFRate:
            plans.extend(self._plan_AdaptiveLIFRate(ops))
        else:
            raise ValueError("Unsupported neuron type '%s'"
                                 % neuron_class.__name__)
    return plans

• in clra_nonlinearities.py

  def plan_adaptive_lif_rate(queue, J, R, ref, tau, dt, tau_adapt, inc, tag=None, n_elements=0):
  assert R.ctype == J.ctype
  
  inputs = dict(j=J)
  outputs = dict(r=R)
  parameters = dict(tau=tau, ref=ref,tau_adapt = tau_adapt, inc = inc)
  textconf = dict(Rtype=R.ctype)
  declares = """
      const ${Rtype} c0 = 0, c1 = 1;
      """
  text = """
      j = max(j - inc, c0);
      r = c1 / (ref + tau * log1p(c1/j));
      inc += (dt/tau_adapt)*(inc*(r-n))
      """
  declares = as_ascii(
      Template(declares, output_encoding='ascii').render(**textconf))
  return _plan_template(
      queue, "cl_adaptive_lif_rate", text, declares=declares,
      tag=tag, n_elements=n_elements,
      inputs=inputs, outputs=outputs, parameters=parameters)

Thanks for your time and your tips :D

Menelik.

Rate mode represents a different neuron type (e.g. a different kernel responsible from mapping input current to output level).

This means adding a new kind of population in Model:

• Implement a small model with rate-mode neurons to test things out.
• Add self.rate_populations = [] in the constructor
• Add a RatePopulation type alongside Population in base.py
• Add alloc_rate_populations and do_rate_populations in npy simulator.

CPU and GPU may in general benefit from different memory layouts. How and when should memory layout decisions be made?

The "decoders" plan created in sim_ocl.py currently makes an inefficient use of a GPU. Each of a relatively small number of signals is decoded by a long inner product. This runs pretty much sequentially right now, but a scan prefix sum thing could be a lot faster.

TODO: write a kernel for ragged gemv that is launched with global size (workgroup_size, n_signals) and which performs each of the A_js inner products with a parallel prefix scan.

Hint: start by reading about parallel prefix sum, and see how it's implemented in the PyOpenCL helper utilities.

(Until I can add you from the dropdown box, I'll just mention you, @zero-impact)

Subtask: you'll have to set up some unit tests for the ragged gemv code to ensure that all versions are consistent with each other and correct. We can discuss how to do that, I'm not imagining anything very sophisticated.

Related: http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html
Related: http://documen.tician.de/pyopencl/algorithm.html#module-pyopencl.scan

https://github.com/ctn-waterloo/nengo_theano/blob/master/nengo/scripts/cortical_action.py

OpenCL offers the possibility of using e.g. float16 or even float8 numbers. Can the simulator still work at such low precision? Does it go faster?

Mako provides a more readable and more expressive template language than Python's native string substituion mechanism.

TODO: refactor the kernel-generators we have to use Mako, and start using Mako for new kernels.

The new nengo.LIF uses a better integration scheme. Some tests are failing right now because we're still using the older, less accurate scheme. (See nengo LIF tests and test_clra_nonlinearities.py)

http://devgurus.amd.com/thread/159457

Reikna is a library for generating higher-level OpenCL implementations. It is similar in spirit to the plan_XXX functions defined in the ocl/ subdirectory. What might be gained by introducing a dependency on Reikna and using the classes etc. in that project?

http://reikna.publicfields.net

https://github.com/ctn-waterloo/nengo_theano/blob/master/nengo/scripts/dot_product.py

Serialization would be a good thing to have, even immediately to make performance tuning easier.

I'm thinking: save all data as numpy arrays, and some sufficient summary of the various graphs to quickly rebuild kernels for whatever OCL device is used at load-time.

nengo.tests.test_neurons.test_lif fails on my Intel chipset due to lack of memory, which is strange because it shouldn't use that much. I think this might have to do with the LIF kernel not being vectorized on this chipset (export PYOPENCL_COMPILER_OUTPUT=1 when running the test for details).

Test the utility of the MultiEnsemble design by implementing the basal ganglia model described here:

https://github.com/ctn-waterloo/nengo_theano/blob/master/nengo/templates/basalganglia.py

Test the opencl implementation as this file tests the nengo_theano one:

https://github.com/ctn-waterloo/nengo_theano/blob/master/nengo/test/test_basal_ganglia.py

Got this error running the latest commit on the master branch (9d4ff91):

Traceback (most recent call last):
  File "run_spaun.py", line 203, in <module>
    sim = nengo_ocl.sim_ocl.Simulator(model, dt=cfg.sim_dt, context=ctx)
  File "/home/abruser/git/nengo_ocl/nengo_ocl/sim_ocl.py", line 65, in __init__
    self, network=network, dt=dt, seed=seed, model=model)
  File "/home/abruser/git/nengo_ocl/nengo_ocl/sim_npy.py", line 429, in __init__
    self._plan.extend(self.plan_op_group(op_type, op_list))
  File "/home/abruser/git/nengo_ocl/nengo_ocl/sim_npy.py", line 439, in plan_op_group
    return getattr(self, 'plan_' + op_type.__name__)(ops)
  File "/home/abruser/git/nengo_ocl/nengo_ocl/sim_ocl.py", line 266, in plan_SimSynapse
    A = self.RaggedArray([f.den for f in steps])
AttributeError: 'functools.partial' object has no attribute 'den'

This was with trying to run Spaun 2.0 with opencl. I was previously on the processes branch and everything built fine there.

Test the MultiEnsemble abstraction by implementing a cleanup memory, as described here:

https://github.com/ctn-waterloo/nengo_theano/blob/master/nengo/test/test_cleanup.py

@hunse Did you investigate Clyther for doing the ast_conversion type stuff? What do you think of it?

http://srossross.github.io/Clyther/index.html

If you see an error about SimDirect not being found, that means you too are out of date with nengo.

Quick fix: downgrade your nengo to commit 53efdc9a1048f983a297146fc536fea1ca077e07

Better fix: patch nengo_ocl to use the new builder stuff.

Right now it computes the wrong answer. Make plan_lif at least recognize the situation and raise an error, or, possibly support at least 2D inputs.

The Numpy simulator in this project is significantly slower than the Numpy reference simulator in Nengo (~4x slower, even using longer simulation times which should diminish any intensive planning tasks that have to be done on the first simulation step). We should see why this is. Also of note is that the OCL simulator on both cpu and gpu seems reasonably fast, so the slowdown doesn't seem to be affecting it (unless it's just blisteringly fast to begin with).