19 Feb
2011
19 Feb
'11
10:49 a.m.
Dear mailing list recipients,
I'm currently working on my bachelor project where I am to do heavy
calculations that are very suited for parallel computation (quantum field
theory using path integrals and markov chain montecarlo). This will be done
with OpenCL in Python, thus using PyOpenCL. This is a quite problematic
approach though since I'm new to both OpenCL and Python. Why? Because
PyOpenCL is a shortcut that blocks me from directly applying 99% of all the
material/tutorials online, since they follow implementation of OpenCL in a C
environment. For every thing I learn about OpenCL, I have to "translate" it
to PyOpenCL. It is possible with this approach, however it feels like a very
slow progress that takes a lot of effort. Since I can't seem to find any
tutorials on learning OpenCL through PyOpenCL on the web, I wonder if anyone
on the mailing list knows of such tutorials? Learning Python isn't the issue
here, and the approach isn't up for change. The question is HOW I'm going to
learn/do this approach in an efficient way.
Thank you for your time!
Yours sincerely,
Patric
20 Feb
20 Feb
3:52 a.m.
Patric Holmvall <patric.hol@...> writes:
...I'm new to both OpenCL and Python ... PyOpenCL
is a shortcut that blocks me
from directly applying 99% of all the material/tutorials online, since they
follow implementation of OpenCL in a C environment.
When I started learning OpenCL I was also new to PyOpenCL, but experienced with
Python itself. As an experienced Python programmer (and having used Numpy a lot
in other projects) it was extremely easy for me to adapt the C examples to
PyOpenCL using the examples as a guide.
But I can see how someone who is new to both of them at the same time would
think that it's a blocker. If I were you I'd start my kernels in C so that you
can deal with programming errors and learning curve stuff without suspecting
your programming environment.
Once you get tired of waiting on recompiles and you understand what you are
doing with these 5 basic things:
1) Input buffers copied from the host
2) Kernel compilation
3) Kernel Argument setup
4) Kernel execution
5) Fetching result buffers back to the host
then I'd learn how to do all that from Python. Once you know enough OpenCL so
that it doesn't look foreign, you will realize that PyOpenCL is not blocking you
at all. It actually makes it more pleasant.
Just my two cents...
--Keith Brafford
5:37 a.m.
Bonjour Patric,
On 19 February 2011 11:49, Patric Holmvall <patric.hol(a)gmail.com> wrote:
Dear mailing list recipients,
I'm currently working on my bachelor project where I am to do heavy
calculations that are very suited for parallel computation (quantum field
theory using path integrals and markov chain montecarlo). This will be done
with OpenCL in Python, thus using PyOpenCL. This is a quite problematic
approach though since I'm new to both OpenCL and Python.
PyOpenCL is the perfect environment to learn OpenCL programming, but
I'm not aware of a step-by-step tutorial based on pyopencl.
You should attack the problem orthogonally ("divide et impera"):
- learn python first. Python is a pleasure to learn and use. The web
is full of very complete and easy to follow tutorials;
- C (for OpenCL) is a bit more difficult to master. C means pointer
and pointer means trouble. Unfortunately you cannot bypass the C
pointers mechanics because they are intensivey used to exchange data
between the two environment. To make the situation worse, opencl
mechanics is not easy to understand because you need to understand the
SIMT (single instruction, multiple threads) computing model first
applied to GPU architecures. I have spent two weeks just to figure out
why my perfectly running code was so slow. NVIDIA and ATI/AMD
implementations are not perfectly equivalent, so you should learn also
how to avoid some pitfalls.
The best tutorial that I have found on the web is this one:
http://developer.amd.com/zones/OpenCLZone/universities/Pages/default.aspx
This is a very generic tutorial that cover AMD/ATI/CELL architectures.
But it is not enough: you will be obliged to read the OpenCL manuals
put on line by NVIDIA and AMD/ATI. This lecture is a must also because
not all the OpenCL devices are made equals: you need to understand
some gory details about hardware ("no guts, no glory").
The biggest bifurcation point is: "Do you need double precision
support (64 bits) _or_ single precision (32 bits) is enough for your
application?".
I'm doing most of my code prototyping using single precision (float)
but, at the end, I will be obliged to switch to double precision in
order to show coherency between the new simulations running on GPU and
the old one running on CPU.
Why? Because
PyOpenCL is a shortcut that blocks me from directly applying 99% of all the
material/tutorials online, since they follow implementation of OpenCL in a C
environment. For every thing I learn about OpenCL, I have to "translate" it
to PyOpenCL. It is possible with this approach, however it feels like a very
slow progress that takes a lot of effort. Since I can't seem to find any
tutorials on learning OpenCL through PyOpenCL on the web, I wonder if anyone
on the mailing list knows of such tutorials? Learning Python isn't the issue
here, and the approach isn't up for change. The question is HOW I'm going to
learn/do this approach in an efficient way.
Finally, (for the moment) GPU programming (also using pyopencl) is
like driving an F1 race car:
you need to invest a lot of time in test driving and tuning. Most of
the thing that you know about "car" does not work on a F1.
A question for opencl development team:
"Is it possible to have the right to commit some - # commented -
example code inside the project?".
Another possibility is to create a completely separate code repository
for the "examples": a lot of people like me will be able to contribute
with examples / application code without take the risk to commit
garbage inside the "core" C/C++/Python engine.
Thanks to all to keep running this fantastic Open Source project :-)
Simone Mannori
11:48 a.m.
Thank you a lot for your answers Simone and Keith!
A separate repository for examples would be very welcome. In fact, if
I could get access to a large set of examples now somehow, I would be
delighted. Especially examples based for AMD/ATI hardware. I find that
I learn more when working with examples rigth now.
A few quick basic questions:
1. If I want to operate on a vector and assign the result to another
vector this is done beautifully in OpenCL, example: fx[gid] =
native_exp(-x[gid]). If I want to sum all the elements in the
resulting vector (fx) to a scalar, is this typically done on the GPU
in OpenCL too (changing it to: fx += native_exp(-x[gid]), or is this
something that is quicker/better to do on the CPU (eg.
numpy.sum(Array))? Is it even possible to sum all the calculations of
the independent work items to the same variable/element?
2. What I'm currently struggling with is how I shall incorporate
workgroup size/local size, thread strides etc and how this is governed
by a) amount of calculations and b) my hardware. Can someone give me
pointers on how I find optimum workgroup size, threads and so on for
Radeon HD5000 series (HD5850, Stream Processors: 1440) and AMD Phenom
II X4 (965BE, 3.4 GHz quadcore) for example 10^9 calculations (10^9
dimensional vectors that I operate on)?
(Or are those parameters governed by the hardware at all?)
Yours sincerely,
Patric
On Sun, Feb 20, 2011 at 6:37 AM, Simone Mannori
<simone.mannori(a)gmail.com> wrote:
Bonjour Patric,
On 19 February 2011 11:49, Patric Holmvall <patric.hol(a)gmail.com> wrote:
Dear mailing list recipients,
I'm currently working on my bachelor project where I am to do heavy
calculations that are very suited for parallel computation (quantum field
theory using path integrals and markov chain montecarlo). This will be done
with OpenCL in Python, thus using PyOpenCL. This is a quite problematic
approach though since I'm new to both OpenCL and Python.
PyOpenCL is the perfect environment to learn OpenCL programming, but
I'm not aware of a step-by-step tutorial based on pyopencl.
You should attack the problem orthogonally ("divide et impera"):
- learn python first. Python is a pleasure to learn and use. The web
is full of very complete and easy to follow tutorials;
- C (for OpenCL) is a bit more difficult to master. C means pointer
and pointer means trouble. Unfortunately you cannot bypass the C
pointers mechanics because they are intensivey used to exchange data
between the two environment. To make the situation worse, opencl
mechanics is not easy to understand because you need to understand the
SIMT (single instruction, multiple threads) computing model first
applied to GPU architecures. I have spent two weeks just to figure out
why my perfectly running code was so slow. NVIDIA and ATI/AMD
implementations are not perfectly equivalent, so you should learn also
how to avoid some pitfalls.
The best tutorial that I have found on the web is this one:
http://developer.amd.com/zones/OpenCLZone/universities/Pages/default.aspx
This is a very generic tutorial that cover AMD/ATI/CELL architectures.
But it is not enough: you will be obliged to read the OpenCL manuals
put on line by NVIDIA and AMD/ATI. This lecture is a must also because
not all the OpenCL devices are made equals: you need to understand
some gory details about hardware ("no guts, no glory").
The biggest bifurcation point is: "Do you need double precision
support (64 bits) _or_ single precision (32 bits) is enough for your
application?".
I'm doing most of my code prototyping using single precision (float)
but, at the end, I will be obliged to switch to double precision in
order to show coherency between the new simulations running on GPU and
the old one running on CPU.
Why? Because
PyOpenCL is a shortcut that blocks me from directly applying 99% of all the
material/tutorials online, since they follow implementation of OpenCL in a C
environment. For every thing I learn about OpenCL, I have to "translate" it
to PyOpenCL. It is possible with this approach, however it feels like a very
slow progress that takes a lot of effort. Since I can't seem to find any
tutorials on learning OpenCL through PyOpenCL on the web, I wonder if anyone
on the mailing list knows of such tutorials? Learning Python isn't the issue
here, and the approach isn't up for change. The question is HOW I'm going to
learn/do this approach in an efficient way.
Finally, (for the moment) GPU programming (also using pyopencl) is
like driving an F1 race car:
you need to invest a lot of time in test driving and tuning. Most of
the thing that you know about "car" does not work on a F1.
A question for opencl development team:
"Is it possible to have the right to commit some - # commented -
example code inside the project?".
Another possibility is to create a completely separate code repository
for the "examples": a lot of people like me will be able to contribute
with examples / application code without take the risk to commit
garbage inside the "core" C/C++/Python engine.
Thanks to all to keep running this fantastic Open Source project :-)
Simone Mannori
1:20 p.m.
Bonjour Patric,
A few quick basic questions:
1. If I want to operate on a vector and assign the result to another
vector this is done beautifully in OpenCL, example: fx[gid] =
native_exp(-x[gid]). If I want to sum all the elements in the
resulting vector (fx) to a scalar, is this typically done on the GPU
in OpenCL too (changing it to: fx += native_exp(-x[gid]), or is this
something that is quicker/better to do on the CPU (eg.
numpy.sum(Array))?
fx += native_exp(-x[gid]) ; //** looks good :-)
Is it even possible to sum all the calculations of
the independent work items to the same variable/element?
I don't know. Wait the answer of more competent people.
2. What I'm currently struggling with is how I shall incorporate
workgroup size/local size, thread strides etc and how this is governed
by a) amount of calculations and b) my hardware. Can someone give me
pointers on how I find optimum workgroup size, threads and so on for
Radeon HD5000 series (HD5850, Stream Processors: 1440) and AMD Phenom
II X4 (965BE, 3.4 GHz quadcore) for example 10^9 calculations (10^9
dimensional vectors that I operate on)?
(Or are those parameters governed by the hardware at all?)
For AMD/ATI hardware you should create workgroup working with 16 items
an regroup it by four (64 work items or threads). NVIDIA is happy too
with workgroup with granularity equal to 64.
You may regroup up to four float (float4) and do the operation in SSE
way (up to four multiply and accumulate in a single clock cycle).
You must read the AMD and NVDIA OpenCL manual. Write simple example
and try to correlate the measured performances with the theorical one.
I'm still on the learning curve, so handle my suggestions with care :-)
Simone
2:52 p.m.
Hi Patric, Keith, Simone, all,
first off, I just wanted to say that I love this community. :)
On Sun, 20 Feb 2011 12:48:27 +0100, Patric Holmvall <patric.hol(a)gmail.com> wrote:
Thank you a lot for your answers Simone and Keith!
A separate repository for examples would be very welcome. In fact, if
I could get access to a large set of examples now somehow, I would be
delighted. Especially examples based for AMD/ATI hardware. I find that
I learn more when working with examples rigth now.
What we have is in the examples/ directory within PyOpenCL. PyCUDA, on
account of having been around longer, has a bunch more examples, here:
http://wiki.tiker.net/PyCuda/Examples
If you think that might be helpful to you, you can try and translate a
few of those to PyOpenCL as a start.
Is it even possible to sum all the calculations of
the independent work items to the same variable/element?
When done in parallel, that falls under the term 'reduction'
[1]. PyOpenCL has code to help you do this [2]. This is a tiny bit
harder to write than vector-vector stuff, but not much. See the source
[3].
[1]
https://secure.wikimedia.org/wikipedia/en/wiki/Fold_(higher-order_function)
[2]
http://documen.tician.de/pyopencl/array.html#module-pyopencl.reduction
[3] http://git.tiker.net/pyopencl.git/blob/HEAD:/pyopencl/reduction.py
Hope this helps,
Andreas
4:40 p.m.
Patric Holmvall <patric.hol@...> writes:
Can someone give me
pointers on how I find optimum workgroup size, threads and so on for
Radeon HD5000 series (HD5850, Stream Processors: 1440) and AMD Phenom
II X4 (965BE, 3.4 GHz quadcore) for example 10^9 calculations (10^9
dimensional vectors that I operate on)?
(Or are those parameters governed by the hardware at all?)
You should join the AMD OpenCL forum for those sorts of questions:
http://forums.amd.com/devforum/categories.cfm?catid=390&entercat=y
My quick answer:
1) Your 1440 stream processors are really 288 VLIW processors with 5 slots, 4
for single precision floating point and 1 fancy slot for other stuff:
http://images.bit-tech.net/content_images/2009/09/ati-radeon-hd-5870-archit…
So for floats, you really have 1152 processors, if you can make heavy use of
the float4 data type. Otherwise you only have 288 processors.
2) As far as workgroup size goes, your 288 VLIW processors are divided into 18
"compute units" (the red rectangles that occupy most of the room in this
picture, which shows 20 since it's the 5870):
http://images.bit-tech.net/content_images/2009/09/ati-radeon-hd-5870-archit…
So if you make your workgroup size a multiple of 16 it will fit nicely into the
architecture. But remember, if you are able to use the float4 data type, you
will be able to do 64 pieces of work in 16 work items.
--Keith Brafford
21 Feb
21 Feb
8:21 p.m.
Thank you Andreas and Keith!
I think I just have a hard time getting my head around how I should
set up my workgroup size/local size, thread strides and so on, but
I'll get around it eventually I guess.
Another topic:
The basic examples covers how to do operations on vectors. I'm trying
to increase the dimensions and want to do operations on a matrix/array
of lets say Nx6 elements.
1. Is there anything wrong with the following approach? (syntax? am I
"thinking wrong" when passing arrays like this to the kernel?)
2. How can I split the other dimension between threads and workgroups
in an optimal way (like x[gid_1][gid_2] for example)? Example code
would be very appreachiated.
3. I suspect that I have a syntax error in the OpenCL code when
accessing elements, like x[gid][2] for example. What is the correct
syntax?
(Notice that this might not be the best example as the calculations
could be reduced to vectors, but I will need to learn how to do
similiar stuf eventually.)
#### Python ####
N = 10**6
x = numpy.random.random(N,6).astype(numpy.float32) # Allocate a Nx6
array of random numbers
x_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=x)
fx_buf = cl.Buffer(ctx, mf.WRITE_ONLY, x.nbytes)
prg.multidim_mc(queue, x.shape, None, x_buf, fx_buf, numpy.uint32(N))
#### OpenCL Kernel ####
prg = cl.Program(ctx, """
__kernel void multidim_mc(__global const float *x, __global float
*fx, uint const N)
{
int gid = get_global_id(0);
if(gid >= N)
{
return;
}
fx[gid] = (x[gid][0]-x[gid][3])^2 + (x[gid][1]-x[gid][4])^2 +
(x[gid][2]-x[gid][5])^2;
}
""").build()
Thanks a lot for a great community!
Yours sincerely,
Patric
On Sun, Feb 20, 2011 at 5:40 PM, Keith Brafford
<keith.brafford(a)gmail.com> wrote:
Patric Holmvall <patric.hol@...> writes:
Can someone give me
pointers on how I find optimum workgroup size, threads and so on for
Radeon HD5000 series (HD5850, Stream Processors: 1440) and AMD Phenom
II X4 (965BE, 3.4 GHz quadcore) for example 10^9 calculations (10^9
dimensional vectors that I operate on)?
(Or are those parameters governed by the hardware at all?)
You should join the AMD OpenCL forum for those sorts of questions:
http://forums.amd.com/devforum/categories.cfm?catid=390&entercat=y
My quick answer:
1) Your 1440 stream processors are really 288 VLIW processors with 5 slots, 4
for single precision floating point and 1 fancy slot for other stuff:
http://images.bit-tech.net/content_images/2009/09/ati-radeon-hd-5870-archit…
So for floats, you really have 1152 processors, if you can make heavy use of
the float4 data type. Otherwise you only have 288 processors.
2) As far as workgroup size goes, your 288 VLIW processors are divided into 18
"compute units" (the red rectangles that occupy most of the room in this
picture, which shows 20 since it's the 5870):
http://images.bit-tech.net/content_images/2009/09/ati-radeon-hd-5870-archit…
So if you make your workgroup size a multiple of 16 it will fit nicely into the
architecture. But remember, if you are able to use the float4 data type, you
will be able to do 64 pieces of work in 16 work items.
--Keith Brafford
_______________________________________________
PyOpenCL mailing list
PyOpenCL(a)tiker.net
http://lists.tiker.net/listinfo/pyopencl
24 Feb
24 Feb
4:12 p.m.
Hi Patric,
On Mon, 21 Feb 2011 21:21:11 +0100, Patric Holmvall <patric.hol(a)gmail.com> wrote:
Thank you Andreas and Keith!
I think I just have a hard time getting my head around how I should
set up my workgroup size/local size, thread strides and so on, but
I'll get around it eventually I guess.
Another topic:
The basic examples covers how to do operations on vectors. I'm trying
to increase the dimensions and want to do operations on a matrix/array
of lets say Nx6 elements.
1. Is there anything wrong with the following approach? (syntax? am I
"thinking wrong" when passing arrays like this to the kernel?)
2. How can I split the other dimension between threads and workgroups
in an optimal way (like x[gid_1][gid_2] for example)? Example code
would be very appreachiated.
3. I suspect that I have a syntax error in the OpenCL code when
accessing elements, like x[gid][2] for example. What is the correct
syntax?
Can you please run your code and say what error message you're getting?
Thanks,
Andreas
3654
days inactive
3659
days old
8 comments
4 participants
participants (4)
-
Andreas Kloeckner -
Keith Brafford -
Patric Holmvall -
Simone Mannori