Re: [PyOpenCL] Grid-based (Py)OpenCL frameworks?
by Karl Czajkowski
On Oct 28, CRV§ADER//KY modulated:
> ... Individual instruments
> of the same type may have more or less sample points for their curve
> attribute.
>
OK, I understand that it would be awkward to pack multiple
instruments' worth of attributes into one NDarray. So you'll need at
least 20K OpenCL jobs, each of which evalutes 2M scenarios. The job
will parallelize scenarios for the same instrument.
This is a very different operating regime than I am familar with. I
might run about 5 jobs to process one whole image, or perhaps 50-100
if I am using sub-block decomposition (repeating the 5 jobs in a
Python loop with different NDarrays as inputs).
> Can the instrument attributes fit in local device memory?
>
> Yes, as I said it's 50MB after compressing, more like 300-500MB
> uncompressed.
> I can optimize though.
>
I don't mean node memory in a cluster, but OpenCL "local memory" that
is shared by compute units in a work group. So, the question is about
the attributes for a single instrument. Ideally, you'd want one
leader to fetch the attributes from OpenCL global memory into local
memory, and then let all work group members reuse these attributes.
The scenario values will be distinct for each work item, so those can
be fetched directly from global memory by each work item. Hopefully
you can vectorize these global loads, whether explicitly in your
kernel or just by auto-vectorization in the OpenCL compiler.
> Didn't start writing the code yet... I'm still in an exploration phase.
>
I don't know how complex one instrument calculation is, but it seems
to me that you should try to write kernels for a couple instrument
sub-classes so you can measure their run times. E.g. try the simplest
kernel (in both programming effort and runtime) and then a more costly
one if it still seems viable to proceed.
With a rough idea of the OpenCL run time for one instrument on all 2M
scenarios, you'll have to return to the per-job preparation cost. I
would be concerned that this cost might dominate the total run time,
if the Python interpreter is doing too much work to marshal the OpenCL
input and output buffers and manage file IO.
Karl
3 years, 7 months
Re: [PyOpenCL] Grid-based (Py)OpenCL frameworks?
by Karl Czajkowski
A few questions and comments in-line below...
On Oct 21, CRV§ADER//KY modulated:
> 1. circa 20,000 Python objects of class "Instrument". Each instrument
> is defined by a subclass and a dict of attributes, which will be a
> handful of scalars most of the times, or a couple of MB worth of numpy
> arrays in the worst case.
> Total size: 50MB after compressing the numpy arrays
>
If I understand correctly, each attribute may be a scalar or a 1D
Numpy array of variable length? Does the attribute shape vary by
individual instrument or by instrument sub-class?
> 2. 120 risk factors, each of which is a numpy 1D array of 2 million
> doubles (the risk factor values for each simulation scenario)
> Total size: 1.8GB
>
So, each logical kernel needs as input one 120 scalar row of risk
factors and one instrument attribute set, and it outputs one scalar?
> My calculation happens in two phases:
> 1.
> Simulation: for every one of the 20,000 instruments, calculate the
> instrument value as a function of the instrument scalar settings and a
> subset of the risk factor vectors. There will be different functions
> (kernels) depending on the instrument subclass. The output is always a
> 1D array of 2 million doubles per instrument - or if you prefer, a 2D
> arrray of 20,000 x 2,000,000 doubles. Some instruments require as input
> the output value of other instruments, in a recursive dependency tree.
> Total output size: 300GB
>
Have you already micro-benchmarked any mappings of this to OpenCL? It
seems to me worth checking:
A. K OpenCL jobs of shape (N,), with each worker evaluating one of K
instruments for one of N scenarios.
B. B OpenCL jobs of shape (K/B, N), with each worker evaluating one
of K instruments for one of N scenarios in one of B blocks.
Can the instrument attributes fit in local device memory? If so, this
can easily benefit (A) and may also help in (B) if you can structure
the global shape (K/B, N) into smaller (W, N) workgroups that share a
single instrument...?
These are of course for K instruments in a single class, so they can
share the same kernel. The output should be a K x N array of scalars,
if I understand your problem statement. I'd limit the numbers K and N
for testing, before worrying about further decomposition to fit the
device and driver limits which probably cannot cope with a 20K much
less 2M job shape axis.
I'd test on both GPU and CPU devices, including existing devices in
your cluster. If your cluster isn't the latest generation of CPUs
and/or GPUs, I'd also try to test on newer equipment; there could be
dramatic performance improvements that would allow a much smaller
number of new devices to meet or exceed a large pool of older ones...
> 2.
> Vertical aggregation:
> I calculate the value of circa 150 nodes, each of which is a vector of
> 2 million doubles defined as a weighted sum of the value of up to 8,000
> instruments (with the weights being scalar): node_value = instr_value1
> * k1 + instr_value2 * k2 + ... + instr_valueN * kn
> Each of the 20,000 instruments can contribute to more than one of the
> output 150 nodes.
>
This phase seems trivially parallelizable and vectorizable. You can
almost dismiss it while optimizing the phase 1 work and overall data
transfers.
Karl
3 years, 7 months
Re: [PyOpenCL] Grid-based (Py)OpenCL frameworks?
by Karl Czajkowski
You've only described a K x N processing problem, where you would run
N kernels that each process one row of K values. You haven't
described any cross-communication or data shuffling if there are
multiple such sub-problems, nor approximate amount of work per input
or output data. Are your tasks truly independent? What data
management do you have to do to get your inputs into a parallel or
distributed decomposition?
At one far extreme, a high throughput job manager could be used to
execute a set of independent PyOpenCL programs, each sized to fit on
your OpenCL devices, each processing a different input file containing
a subset of your N rows of data.
In the middle are a huge number of choices to balance IO, memory, and
compute resources. This leads to a huge number of different research
programs all focusing on different niches and machine models.
If you really want to abstract away the GPGPU devices, you might want
to look at OpenMP or similar projects that have tried to add such
devices to their targets for auto-vectorization. I don't work in that
space, and so have no specific recommendations.
At the other extreme, I adopted PyOpenCL to allow me to do my ad-hoc
processing in Python and Numpy with OpenCL callouts for certain
bottlenecks. I have some image processing tasks where there isn't
even enough compute time per byte of input to warrant the PCIe
transfer bottleneck in some cases. It is the same speed to run on an
i7-4700MQ mobile quad-core CPU (using just SIMD+multicore) as to run
on a desktop Kepler GPU.
For me, the data IO from disk or network would also dominate, so
distributed processing is pointless. Even still, I have used
explicit sub-block decomposition to split my large images into smaller
OpenCL tasks that can be marshaled through the system RAM or GPU to
improve locality and limit the intermediate working set sizes.
Karl
3 years, 7 months
Re: [PyOpenCL] Grid-based (Py)OpenCL frameworks?
by Andreas Kloeckner
"CRV§ADER//KY" <crusaderky(a)gmail.com> writes:
> Hi all,
> I'm looking into setting up a cluster of GPGPU nodes. The nodes would be
> Linux based, and communicate between each other via ethernet. Each node
> would have multiple GPUs.
>
> I need to run a problem that for 99% can be described as y[i] = f(x1[i],
> x2[i], ... xn[i]), running on 1D vectors of data. In other words, I have n
> input vectors and 1 output vector, all of the same size, and worker i-th
> will exclusively need to access element i-th of every vector.
>
> Are there any frameworks, preferably in Python and with direct access to
> OpenCL, that allow to transparently split the input data in segments, send
> them over the network, do caching, feeding executor queues, etc. etc.?
>
> Data reuse is very heavy so if a vector is already in VRAM I don't want to
> load it twice.
>
> Also, are there PyOpenCL bolt-ons that allow for virtual VRAM? That is, to
> have more buffers than you can fit in VRAM, and transparently swap to
> system RAM those thare are not immediately needed?
VirtCL is one. There was another, but I forgot what it was called.
https://dl.acm.org/citation.cfm?id=2688505
Andreas
3 years, 7 months
Grid-based (Py)OpenCL frameworks?
by CRV§ADER//KY
Hi all,
I'm looking into setting up a cluster of GPGPU nodes. The nodes would be
Linux based, and communicate between each other via ethernet. Each node
would have multiple GPUs.
I need to run a problem that for 99% can be described as y[i] = f(x1[i],
x2[i], ... xn[i]), running on 1D vectors of data. In other words, I have n
input vectors and 1 output vector, all of the same size, and worker i-th
will exclusively need to access element i-th of every vector.
Are there any frameworks, preferably in Python and with direct access to
OpenCL, that allow to transparently split the input data in segments, send
them over the network, do caching, feeding executor queues, etc. etc.?
Data reuse is very heavy so if a vector is already in VRAM I don't want to
load it twice.
Also, are there PyOpenCL bolt-ons that allow for virtual VRAM? That is, to
have more buffers than you can fit in VRAM, and transparently swap to
system RAM those thare are not immediately needed?
Thanks
Guido
3 years, 7 months
Re: [PyOpenCL] invalid mem object when splitting kernel enqueue and set_args across class methods
by Blair Azzopardi
Looks like I've found a workaround. It involves setting the set_args
parameters as object properties.
def run_set_args(self):
> mf = cl.mem_flags
> self.a_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
> hostbuf=a_np)
> self.b_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
> hostbuf=b_np)
> self.res_g = cl.Buffer(self.ctx, mf.WRITE_ONLY, a_np.nbytes)
> self.sum_k = self.prg.sum
> self.sum_k.set_args(self.a_g, self.b_g, self.res_g)
On Sun, Oct 18, 2015 at 4:13 PM, Blair Azzopardi <blairuk(a)gmail.com> wrote:
> Hi
>
> I'm trying to split my kernel invocations into class methods. Basically,
> I'd like to call set_args in one method and then enqueue_nd_range_kernel in
> another. However, when I try this I receive a "clEnqueueNDRangeKernel
> failed: invalid mem object" exception.
>
> I've adapted the front page example into a short snippet that demonstrates
> the problem. Can anyone see if there's something obvious I', missing. If
> not, I'll file an issue on gitbub.
>
>
>> import numpy as np
>> import pyopencl as cl
>> a_np = np.random.rand(50000).astype(np.float32)
>> b_np = np.random.rand(50000).astype(np.float32)
>> class Test(object):
>> def __init__(self):
>> self.ctx = cl.create_some_context()
>> self.queue = cl.CommandQueue(self.ctx)
>> self.prg = cl.Program(self.ctx, """
>> __kernel void sum(__global const float *a_g, __global const float
>> *b_g, __global float *res_g) {
>> int gid = get_global_id(0);
>> res_g[gid] = a_g[gid] + b_g[gid];
>> }
>> """).build()
>> def run_all(self):
>> mf = cl.mem_flags
>> a_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
>> hostbuf=a_np)
>> b_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
>> hostbuf=b_np)
>> res_g = cl.Buffer(self.ctx, mf.WRITE_ONLY, a_np.nbytes)
>> sum_k = self.prg.sum
>> sum_k.set_args(a_g, b_g, res_g)
>> cl.enqueue_nd_range_kernel(self.queue, sum_k, a_np.shape, None)
>> def run_set_args(self):
>> mf = cl.mem_flags
>> a_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
>> hostbuf=a_np)
>> b_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
>> hostbuf=b_np)
>> res_g = cl.Buffer(self.ctx, mf.WRITE_ONLY, a_np.nbytes)
>> self.sum_k = self.prg.sum
>> self.sum_k.set_args(a_g, b_g, res_g)
>>
>> def run_enqueue(self):
>> cl.enqueue_nd_range_kernel(self.queue, self.sum_k, a_np.shape,
>> None)
>> o = Test()
>> # runs fine
>> o.run_all()
>> # run separately gives "clEnqueueNDRangeKernel failed: invalid mem object"
>> o.run_set_args()
>> o.run_enqueue()
>
>
> Thanks
> Blair
>
3 years, 7 months
invalid mem object when splitting kernel enqueue and set_args across class methods
by Blair Azzopardi
Hi
I'm trying to split my kernel invocations into class methods. Basically,
I'd like to call set_args in one method and then enqueue_nd_range_kernel in
another. However, when I try this I receive a "clEnqueueNDRangeKernel
failed: invalid mem object" exception.
I've adapted the front page example into a short snippet that demonstrates
the problem. Can anyone see if there's something obvious I', missing. If
not, I'll file an issue on gitbub.
> import numpy as np
> import pyopencl as cl
> a_np = np.random.rand(50000).astype(np.float32)
> b_np = np.random.rand(50000).astype(np.float32)
> class Test(object):
> def __init__(self):
> self.ctx = cl.create_some_context()
> self.queue = cl.CommandQueue(self.ctx)
> self.prg = cl.Program(self.ctx, """
> __kernel void sum(__global const float *a_g, __global const float
> *b_g, __global float *res_g) {
> int gid = get_global_id(0);
> res_g[gid] = a_g[gid] + b_g[gid];
> }
> """).build()
> def run_all(self):
> mf = cl.mem_flags
> a_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
> hostbuf=a_np)
> b_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
> hostbuf=b_np)
> res_g = cl.Buffer(self.ctx, mf.WRITE_ONLY, a_np.nbytes)
> sum_k = self.prg.sum
> sum_k.set_args(a_g, b_g, res_g)
> cl.enqueue_nd_range_kernel(self.queue, sum_k, a_np.shape, None)
> def run_set_args(self):
> mf = cl.mem_flags
> a_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
> hostbuf=a_np)
> b_g = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR,
> hostbuf=b_np)
> res_g = cl.Buffer(self.ctx, mf.WRITE_ONLY, a_np.nbytes)
> self.sum_k = self.prg.sum
> self.sum_k.set_args(a_g, b_g, res_g)
>
> def run_enqueue(self):
> cl.enqueue_nd_range_kernel(self.queue, self.sum_k, a_np.shape,
> None)
> o = Test()
> # runs fine
> o.run_all()
> # run separately gives "clEnqueueNDRangeKernel failed: invalid mem object"
> o.run_set_args()
> o.run_enqueue()
Thanks
Blair
3 years, 7 months
Re: [PyOpenCL] PyOpenCL error in python2.7*
by Andreas Kloeckner
Longgang Pang <pang(a)fias.uni-frankfurt.de> writes:
> I have problem when run the most recent PyOpenCL2015 with python2.7.6
> and 2.7.8
> The full output is as following, looks like 'ArgumentError' has no
> attribute 'what' for python version 2.7.6 and 2.7.8?
> There is no problem when I run pyopencl with python/3.3.6 or python/2.7.10.
>
> Do you know how to fix it in the code without modifying pyopencl?
Can you submit some code that reproduces this? ArgumentError should not
be caught by this handler.
Andreas
3 years, 7 months