work group size errors

[Tim Hutton]

You might have seen the announcement of 0.2 I put on conwaylife.com:
http://www.conwaylife.com/forums/viewtopic.php?f=11&t=869

A couple of users reported OpenCL problems:

Wojowu: "Failed to open file. Error:

OpenCL_RD::ReloadKernelIfNeeded : expecting CL_KERNEL_WORK_GROUP_SIZE
to be a power of 2"

and

Hektor: "An error occurred when running the simulation:

OpenCL_RD::Update : clEnqueueNDRangeKernel failed: Invalid work group size"

I think both of these are to do with bad assumptions I was making in
choosing the local work group size - the number of cells that are
processed in parallel. I've just discovered (r722) that we actually
don't need to specify the local work group size at all - we can pass
NULL instead and let the OpenCL system choose. On my system this also
improves the speed by 25% on some patterns. I've also improved the
OpenCL diagnostics function to return more details (r720) which might
have helped with diagnosing these problems.

I'm 80% confident that this will fix the problems reported, so if you
can test it for me on your machines that would be great. I'm
interested in:
a) whether everything still works.
b) whether the speed is faster, slower, or about the same.

If it's all OK then we can deliver a new build.

--
Tim Hutton - http://www.sq3.org.uk - http://profiles.google.com/tim.hutton/

[Robert Munafo]

svn update gave me version r722. It's been a month or two since I tried to build Ready. Here's what I get now:

 : make

-- Configuring done

-- Generating done

-- Build files have been written to: /Users/munafo/shared/proj/hutton-gray-scott/Ready_gui_build

Scanning dependencies of target CopyFiles

[  1%] Generating Patterns/grayscott_demo.vti

[  3%] Generating Patterns/meinhardt_stripes.vti

[  5%] Generating Patterns/schlogl.vti

[  7%] Generating Patterns/heat_equation.vti

[  9%] Generating Patterns/turing.vti

[ 11%] Generating Patterns/kernel_test.vti

[ 12%] Generating Patterns/conway_life.vti

[ 14%] Generating Patterns/CPU-only/grayscott_1D.vti

[ 16%] Generating Patterns/CPU-only/grayscott_2D.vti

[ 18%] Generating Patterns/CPU-only/grayscott_3D.vti

[ 20%] Generating Patterns/FitzHugh-Nagumo/tip-splitting.vti

[ 22%] Generating Patterns/FitzHugh-Nagumo/spiral_turbulence.vti

[ 24%] Generating Patterns/FitzHugh-Nagumo/pulsate.vti

[ 25%] Generating Patterns/FitzHugh-Nagumo/squid_axon.vti

[ 27%] Generating Help/about.gif

[ 29%] Generating Help/about.html

[ 31%] Generating Help/action.html

[ 33%] Generating Help/credits.html

[ 35%] Generating Help/file.html

[ 37%] Generating Help/help.html

[ 38%] Generating Help/mouse.html

[ 40%] Generating Help/quickstart.html

[ 42%] Generating Help/tips.html

[ 44%] Generating Help/changes.html

[ 46%] Generating Help/edit.html

[ 48%] Generating Help/formats.html

[ 50%] Generating Help/index.html

[ 51%] Generating Help/problems.html

[ 53%] Generating Help/view.html

[ 55%] Generating Help/introduction.html

[ 57%] Generating README.txt

[ 59%] Generating COPYING.txt

[ 61%] Generating TODO.txt

[ 62%] Generating BUILD.txt

[ 62%] Built target CopyFiles

Scanning dependencies of target app_bundle

[ 62%] Built target app_bundle

Scanning dependencies of target readybase

[ 64%] Building CXX object CMakeFiles/readybase.dir/BaseRD.cpp.o

[ 66%] Building CXX object CMakeFiles/readybase.dir/GrayScott.cpp.o

[ 68%] Building CXX object CMakeFiles/readybase.dir/OpenCL_RD.cpp.o

[ 70%] Building CXX object CMakeFiles/readybase.dir/OpenCL_utils.cpp.o

/Users/munafo/shared/proj/hutton-gray-scott/reaction-diffusion/Ready/OpenCL_utils.cpp: In function ‘void throwOnError(cl_int, const char*)’:

/Users/munafo/shared/proj/hutton-gray-scott/reaction-diffusion/Ready/OpenCL_utils.cpp:31: error: ‘runtime_error’ was not declared in this scope

/Users/munafo/shared/proj/hutton-gray-scott/reaction-diffusion/Ready/OpenCL_utils.cpp: In function ‘const char* GetDeviceInfoIdAsString(cl_device_info)’:

/Users/munafo/shared/proj/hutton-gray-scott/reaction-diffusion/Ready/OpenCL_utils.cpp:366: error: ‘CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF’ was not declared in this scope

/Users/munafo/shared/proj/hutton-gray-scott/reaction-diffusion/Ready/OpenCL_utils.cpp:367: error: ‘CL_DEVICE_HOST_UNIFIED_MEMORY’ was not declared in this scope

/Users/munafo/shared/proj/hutton-gray-scott/reaction-diffusion/Ready/OpenCL_utils.cpp:368: error: ‘CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR’ was not declared in this scope

(and more of the same)

I'm getting the exact same errors on two different machines (a Nehalem 8-core and a Sandy bridge 4-core) both of which are running Mac OS 10.6.8. One has cmake version 2.8.5 and the other has 2.8.7. I suppose the error has to do with my OpenCL installation, but I don't remember how to ascertain what versions of the OpenCL headers/libraries I have (except that I'm pretty sure it came with Xcode tools)

- Robert

On Mon, Mar 26, 2012 at 12:43, Tim Hutton <tim.h...@gmail.com> wrote:

[...]

I'm 80% confident that this will fix the problems reported, so if you
can test it for me on your machines that would be great. I'm
interested in:
a) whether everything still works.
b) whether the speed is faster, slower, or about the same.

If it's all OK then we can deliver a new build.

--
  Robert Munafo  --  mrob.com
  Follow me at: gplus.to/mrob - fb.com/mrob27 - twitter.com/mrob_27 -
mrob27.wordpress.com - youtube.com/user/mrob143 - rilybot.blogspot.com

[Tim Hutton]

Hi Robert,

The runtime_error issue I've fixed. (I think it's because MSVC's
implementation of STL is a bit different.)

The CL_* things are due to the Apple build using whichever OpenCL.h is
installed instead of our OpenCL_Dyn_Load.h. If you can give me the
full list of errors I'll exclude those. It's just for reporting, so it
doesn't matter if some are missing.

--

[Andrew Trevorrow]

Tim:

> ...

> I'm 80% confident that this will fix the problems reported, so if you
> can test it for me on your machines that would be great. I'm
> interested in:
> a) whether everything still works.

I had to add some defines to OpenCL_utils.hpp that are missing from
Apple's OpenCL headers. (But I just saw your reply to Robert's report
so feel free to remove my additions if you think it's better simply
not to use them in the diagnostics. I've appended my diagnostic
output in case it reveals anything interesting.)

> b) whether the speed is faster, slower, or about the same.

A bit slower unfortunately. Mostly around 10% slower, but pulsate.vti
dropped from 1800 mcgs to 1200 mcgs. Still fast enough for me, so
I'm quite hapopy to live with this if it means Mac users no longer
see any errors about work group size.

Andrew

------------------------

Found 1 platform(s):
Platform 1:
CL_PLATFORM_PROFILE : FULL_PROFILE
CL_PLATFORM_VERSION : OpenCL 1.0 (Jan 2 2011 18:00:11)
CL_PLATFORM_NAME : Apple
CL_PLATFORM_VENDOR : Apple
CL_PLATFORM_EXTENSIONS :

Found 2 device(s) on this platform.
Device 1:
CL_DEVICE_EXTENSIONS : cl_APPLE_gl_sharing
CL_DEVICE_NAME : ATI Radeon HD 6970M
CL_DEVICE_PROFILE : FULL_PROFILE
CL_DEVICE_VENDOR : AMD
CL_DEVICE_VERSION : OpenCL 1.0
CL_DRIVER_VERSION : 1.0
CL_DEVICE_ADDRESS_BITS : 32
CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE : 0
CL_DEVICE_MAX_CLOCK_FREQUENCY : 149
CL_DEVICE_MAX_COMPUTE_UNITS : 10
CL_DEVICE_MAX_CONSTANT_ARGS : 8
CL_DEVICE_MAX_READ_IMAGE_ARGS : 0
CL_DEVICE_MAX_SAMPLERS : 128
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : 3
CL_DEVICE_MAX_WRITE_IMAGE_ARGS : 0
CL_DEVICE_MEM_BASE_ADDR_ALIGN : 32768
CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE : 128
CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT : 4
CL_DEVICE_VENDOR_ID : 16915200
CL_DEVICE_AVAILABLE : yes
CL_DEVICE_COMPILER_AVAILABLE : yes
CL_DEVICE_ENDIAN_LITTLE : no
CL_DEVICE_ERROR_CORRECTION_SUPPORT : no
CL_DEVICE_IMAGE_SUPPORT : no
CL_DEVICE_GLOBAL_MEM_CACHE_SIZE : 0
CL_DEVICE_GLOBAL_MEM_SIZE : 536870912
CL_DEVICE_LOCAL_MEM_SIZE : 32768
CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE : 65536
CL_DEVICE_MAX_MEM_ALLOC_SIZE : 134217728
CL_DEVICE_IMAGE2D_MAX_HEIGHT : 8192
CL_DEVICE_IMAGE2D_MAX_WIDTH : 8192
CL_DEVICE_IMAGE3D_MAX_DEPTH : 0
CL_DEVICE_IMAGE3D_MAX_HEIGHT : 0
CL_DEVICE_IMAGE3D_MAX_WIDTH : 0
CL_DEVICE_MAX_PARAMETER_SIZE : 1024
CL_DEVICE_MAX_WORK_GROUP_SIZE : 1024
CL_DEVICE_PROFILING_TIMER_RESOLUTION : 37
CL_DEVICE_MAX_WORK_ITEM_SIZES : 1024, 1024, 1024
CL_DEVICE_TYPE : GPU
Device 2:
CL_DEVICE_EXTENSIONS : cl_khr_fp64 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_APPLE_gl_sharing cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions
CL_DEVICE_NAME : Intel(R) Core(TM) i5-2400 CPU @ 3.10GHz
CL_DEVICE_PROFILE : FULL_PROFILE
CL_DEVICE_VENDOR : Intel
CL_DEVICE_VERSION : OpenCL 1.0
CL_DRIVER_VERSION : 1.0
CL_DEVICE_ADDRESS_BITS : 64
CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE : 64
CL_DEVICE_MAX_CLOCK_FREQUENCY : 3100
CL_DEVICE_MAX_COMPUTE_UNITS : 4
CL_DEVICE_MAX_CONSTANT_ARGS : 8
CL_DEVICE_MAX_READ_IMAGE_ARGS : 128
CL_DEVICE_MAX_SAMPLERS : 16
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : 3
CL_DEVICE_MAX_WRITE_IMAGE_ARGS : 8
CL_DEVICE_MEM_BASE_ADDR_ALIGN : 1024
CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE : 128
CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT : 4
CL_DEVICE_VENDOR_ID : 16909312
CL_DEVICE_AVAILABLE : yes
CL_DEVICE_COMPILER_AVAILABLE : yes
CL_DEVICE_ENDIAN_LITTLE : yes
CL_DEVICE_ERROR_CORRECTION_SUPPORT : no
CL_DEVICE_IMAGE_SUPPORT : yes
CL_DEVICE_GLOBAL_MEM_CACHE_SIZE : 6291456
CL_DEVICE_GLOBAL_MEM_SIZE : 12884901888
CL_DEVICE_LOCAL_MEM_SIZE : 16384
CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE : 65536
CL_DEVICE_MAX_MEM_ALLOC_SIZE : 3221225472
CL_DEVICE_IMAGE2D_MAX_HEIGHT : 8192
CL_DEVICE_IMAGE2D_MAX_WIDTH : 8192
CL_DEVICE_IMAGE3D_MAX_DEPTH : 2048
CL_DEVICE_IMAGE3D_MAX_HEIGHT : 2048
CL_DEVICE_IMAGE3D_MAX_WIDTH : 2048
CL_DEVICE_MAX_PARAMETER_SIZE : 4096
CL_DEVICE_MAX_WORK_GROUP_SIZE : 1
CL_DEVICE_PROFILING_TIMER_RESOLUTION : 1
CL_DEVICE_MAX_WORK_ITEM_SIZES : 1, 1, 1
CL_DEVICE_TYPE : CPU

[Robert Munafo]

With revision 725, build works now, thank you Andrew!

Now on to what you actually wanted to know:

On Mon, Mar 26, 2012 at 18:59, Andrew Trevorrow <and...@trevorrow.com> wrote:

> a) whether everything still works.

What should I try (in other words, define "everything" :-) ?

I clicked on all the patterns, and hit Space, Run (return key), Reset (command-R), Step By N (tab), and Run Slower (-) a few times, and it did what I expected.

 

> b) whether the speed is faster, slower, or about the same.

How can I measure the speed (which pattern do I run, how many timesteps per N, and how big (in pixels) should the main view pane be?) 

More importantly, how do I get and build an older version to compare it to? Do I need to delete my entire repository or is there a way to tell SVN to send me back in time? Which revision should I ask it to give me?

[Andrew Trevorrow]

Robert:

> What should I try (in other words, define "everything" :-) ?
>
> I clicked on all the patterns, and hit Space, Run (return key), Reset
> (command-R), Step By N (tab), and Run Slower (-) a few times, and it
> did what I expected.

That sounds fine. I'd be curious to see the output from
Action > Show OpenCL Diagnostics.

> How can I measure the speed ...

> More importantly, how do I get and build an older version to compare it to?

Just download Ready-0.2.1-Mac.zip from the Ready website and compare
its speed with your latest build. Run a few patterns for a few secs
and compare the status bar numbers showing fps and mcgs.

Andrew

[Robert Munafo]

Okay, here's the 8-core Nehalem workstation, with an ATI Radeon HD 5770 GPU:

Found 1 platform(s):

Platform 1:

CL_PLATFORM_PROFILE : FULL_PROFILE

CL_PLATFORM_VERSION : OpenCL 1.0 (Dec 23 2010 17:30:26)

CL_PLATFORM_NAME : Apple

CL_PLATFORM_VENDOR : Apple

CL_PLATFORM_EXTENSIONS : 

Found 3 device(s) on this platform.

Device 1:

CL_DEVICE_EXTENSIONS : cl_APPLE_gl_sharing

CL_DEVICE_NAME : ATI Radeon HD 5770

CL_DEVICE_PROFILE : FULL_PROFILE

CL_DEVICE_VENDOR : AMD

CL_DEVICE_VERSION : OpenCL 1.0 

CL_DRIVER_VERSION : 1.0

CL_DEVICE_ADDRESS_BITS : 32

CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE : 0

CL_DEVICE_MAX_CLOCK_FREQUENCY : 1195

CL_DEVICE_EXTENSIONS : cl_khr_byte_addressable_store cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_APPLE_gl_sharing cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions 

CL_DEVICE_NAME : GeForce GT 120

CL_DEVICE_PROFILE : FULL_PROFILE

CL_DEVICE_VENDOR : NVIDIA

CL_DEVICE_VERSION : OpenCL 1.0 

CL_DRIVER_VERSION : CLH 1.0

CL_DEVICE_ADDRESS_BITS : 32

CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE : 0

CL_DEVICE_MAX_CLOCK_FREQUENCY : 1400

CL_DEVICE_MAX_COMPUTE_UNITS : 4

CL_DEVICE_MAX_CONSTANT_ARGS : 9

CL_DEVICE_MAX_READ_IMAGE_ARGS : 128

CL_DEVICE_MAX_SAMPLERS : 16

CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : 3

CL_DEVICE_MAX_WRITE_IMAGE_ARGS : 8

CL_DEVICE_MEM_BASE_ADDR_ALIGN : 1024

CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE : 128

CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT : 1

CL_DEVICE_VENDOR_ID : 16918016

CL_DEVICE_AVAILABLE : yes

CL_DEVICE_COMPILER_AVAILABLE : yes

CL_DEVICE_ENDIAN_LITTLE : yes

CL_DEVICE_ERROR_CORRECTION_SUPPORT : no

CL_DEVICE_IMAGE_SUPPORT : yes

CL_DEVICE_GLOBAL_MEM_CACHE_SIZE : 0

CL_DEVICE_GLOBAL_MEM_SIZE : 536870912

CL_DEVICE_LOCAL_MEM_SIZE : 16384

CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE : 65536

CL_DEVICE_MAX_MEM_ALLOC_SIZE : 134217728

CL_DEVICE_IMAGE2D_MAX_HEIGHT : 4096

CL_DEVICE_IMAGE2D_MAX_WIDTH : 4096

CL_DEVICE_IMAGE3D_MAX_DEPTH : 2048

CL_DEVICE_IMAGE3D_MAX_HEIGHT : 2048

CL_DEVICE_IMAGE3D_MAX_WIDTH : 2048

CL_DEVICE_MAX_PARAMETER_SIZE : 4352

CL_DEVICE_MAX_WORK_GROUP_SIZE : 512

CL_DEVICE_PROFILING_TIMER_RESOLUTION : 1000

CL_DEVICE_MAX_WORK_ITEM_SIZES : 512, 512, 64

CL_DEVICE_TYPE : GPU

Device 3:

CL_DEVICE_EXTENSIONS : cl_khr_fp64 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_APPLE_gl_sharing cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions 

CL_DEVICE_NAME : Intel(R) Xeon(R) CPU           E5520  @ 2.27GHz

CL_DEVICE_PROFILE : FULL_PROFILE

CL_DEVICE_VENDOR : Intel

CL_DEVICE_VERSION : OpenCL 1.0 

CL_DRIVER_VERSION : 1.0

CL_DEVICE_ADDRESS_BITS : 64

CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE : 64

CL_DEVICE_MAX_CLOCK_FREQUENCY : 2260

CL_DEVICE_MAX_COMPUTE_UNITS : 16

CL_DEVICE_MAX_CONSTANT_ARGS : 8

CL_DEVICE_MAX_READ_IMAGE_ARGS : 128

CL_DEVICE_MAX_SAMPLERS : 16

CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : 3

CL_DEVICE_MAX_WRITE_IMAGE_ARGS : 8

CL_DEVICE_MEM_BASE_ADDR_ALIGN : 1024

CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE : 128

CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT : 4

CL_DEVICE_VENDOR_ID : 16909312

CL_DEVICE_AVAILABLE : yes

CL_DEVICE_COMPILER_AVAILABLE : yes

CL_DEVICE_ENDIAN_LITTLE : yes

CL_DEVICE_ERROR_CORRECTION_SUPPORT : no

CL_DEVICE_IMAGE_SUPPORT : yes

CL_DEVICE_GLOBAL_MEM_CACHE_SIZE : 8388608

CL_DEVICE_GLOBAL_MEM_SIZE : 20937965568

CL_DEVICE_LOCAL_MEM_SIZE : 16384

CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE : 65536

CL_DEVICE_MAX_MEM_ALLOC_SIZE : 5234491392

CL_DEVICE_IMAGE2D_MAX_HEIGHT : 8192

CL_DEVICE_IMAGE2D_MAX_WIDTH : 8192

CL_DEVICE_IMAGE3D_MAX_DEPTH : 2048

CL_DEVICE_IMAGE3D_MAX_HEIGHT : 2048

CL_DEVICE_IMAGE3D_MAX_WIDTH : 2048

CL_DEVICE_MAX_PARAMETER_SIZE : 4096

CL_DEVICE_MAX_WORK_GROUP_SIZE : 1

CL_DEVICE_PROFILING_TIMER_RESOLUTION : 1

CL_DEVICE_MAX_WORK_ITEM_SIZES : 1, 1, 1

CL_DEVICE_TYPE : CPU

The download Andrew is referring to is here:

  http://code.google.com/p/reaction-diffusion/downloads/detail?name=Ready-0.2.1-Mac.zip#makechanges

For some of the patterns, the mcgs figures jump around a lot, so I doubt their accuracy. For example, if I select conway_life.vti and hit Return twice (to start and stop), I can see a "snapshot" of the mcgs figure; out of 20 consecutive samples I got a low of 88 mcgs and a high of 407 mcgs; 6 of the samples were below 100 and the other 14 samples were above 300.

Also, all 3 of the CPU-only ones always say 7 mcgs.

Only the FitzHugh-Nagumo and turing.vti were stable enough for me to consider the numbers worth reporting.

pulsate.vti: Ready-0.2.1-Mac ran at 984 mcgs; current runs at 865 mcgs

spiral_turbulence.vti: Ready-0.2.1-Mac ran at 517 mcgs; current runs at 474 mcgs

squid_axon: always reports 1 mcgs

tip-splitting.vti: numbers jump around too much

turing.vti: Ready-0.2.1-Mac ran at  169 mcgs; current runs at 151 mcgs

Actually, now it looks like ALL of the mcgs numbers jump around a lot. I think we need to apply a decaying-average filter function to that calculation.

I'll email separately with results from the Sandy Bridge system.

- Robert

On Mon, Mar 26, 2012 at 21:13, Andrew Trevorrow <and...@trevorrow.com> wrote:

That sounds fine.  I'd be curious to see the output from
Action > Show OpenCL Diagnostics.

Just download Ready-0.2.1-Mac.zip from the Ready website and compare its speed with your latest build.  Run a few patterns for a few secs and compare the status bar numbers showing fps and mcgs.

[Robert Munafo]

On the Sandy Bridge system, I cannot run any patterns except the three CPU-only ones. The rest give me this error:

Failed to open file. Error:

OpenCL_RD::ReloadContextIfNeeded : Failed to create command queue: Invalid value

Using a printf I discovered that num_devices == 1, but haven't debugged it further.

Here are the diagnostics:

Found 1 platform(s):

Platform 1:

CL_PLATFORM_PROFILE : FULL_PROFILE

CL_PLATFORM_VERSION : OpenCL 1.0 (Dec 26 2010 12:52:21)

CL_PLATFORM_NAME : Apple

CL_PLATFORM_VENDOR : Apple

CL_PLATFORM_EXTENSIONS : 

Found 2 device(s) on this platform.

Device 1:

CL_DEVICE_EXTENSIONS : cl_APPLE_gl_sharing

CL_DEVICE_NAME : ATI Radeon HD 6750M

CL_DEVICE_PROFILE : FULL_PROFILE

CL_DEVICE_VENDOR : AMD

CL_DEVICE_VERSION : OpenCL 1.0 

CL_DRIVER_VERSION : 1.0

CL_DEVICE_ADDRESS_BITS : 32

CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE : 0

CL_DEVICE_MAX_CLOCK_FREQUENCY : 150

CL_DEVICE_MAX_COMPUTE_UNITS : 5

CL_DEVICE_EXTENSIONS : cl_khr_fp64 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_APPLE_gl_sharing cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions 

CL_DEVICE_NAME : Intel(R) Core(TM) i7-2720QM CPU @ 2.20GHz

CL_DEVICE_PROFILE : FULL_PROFILE

CL_DEVICE_VENDOR : Intel

CL_DEVICE_VERSION : OpenCL 1.0 

CL_DRIVER_VERSION : 1.0

CL_DEVICE_ADDRESS_BITS : 64

CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE : 64

CL_DEVICE_MAX_CLOCK_FREQUENCY : 2200

CL_DEVICE_MAX_COMPUTE_UNITS : 8

CL_DEVICE_MAX_CONSTANT_ARGS : 8

CL_DEVICE_MAX_READ_IMAGE_ARGS : 128

CL_DEVICE_MAX_SAMPLERS : 16

CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : 3

CL_DEVICE_MAX_WRITE_IMAGE_ARGS : 8

CL_DEVICE_MEM_BASE_ADDR_ALIGN : 1024

CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE : 128

CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT : 4

CL_DEVICE_VENDOR_ID : 16909312

CL_DEVICE_AVAILABLE : yes

CL_DEVICE_COMPILER_AVAILABLE : yes

CL_DEVICE_ENDIAN_LITTLE : yes

CL_DEVICE_ERROR_CORRECTION_SUPPORT : no

CL_DEVICE_IMAGE_SUPPORT : yes

CL_DEVICE_GLOBAL_MEM_CACHE_SIZE : 6291456

CL_DEVICE_GLOBAL_MEM_SIZE : 3221225472

CL_DEVICE_LOCAL_MEM_SIZE : 16384

CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE : 65536

CL_DEVICE_MAX_MEM_ALLOC_SIZE : 1073741824

CL_DEVICE_IMAGE2D_MAX_HEIGHT : 8192

CL_DEVICE_IMAGE2D_MAX_WIDTH : 8192

CL_DEVICE_IMAGE3D_MAX_DEPTH : 2048

CL_DEVICE_IMAGE3D_MAX_HEIGHT : 2048

CL_DEVICE_IMAGE3D_MAX_WIDTH : 2048

CL_DEVICE_MAX_PARAMETER_SIZE : 4096

CL_DEVICE_MAX_WORK_GROUP_SIZE : 1

CL_DEVICE_PROFILING_TIMER_RESOLUTION : 1

CL_DEVICE_MAX_WORK_ITEM_SIZES : 1, 1, 1

CL_DEVICE_TYPE : CPU

[Tim Hutton]

On 27 March 2012 01:27, Robert Munafo <mro...@gmail.com> wrote:
> With revision 725, build works now, thank you Andrew!
>
> Now on to what you actually wanted to know:
>
> On Mon, Mar 26, 2012 at 18:59, Andrew Trevorrow <and...@trevorrow.com>
> wrote:
>>
>> > a) whether everything still works.
>
>
> What should I try (in other words, define "everything" :-) ?
>
> I clicked on all the patterns, and hit Space, Run (return key), Reset
> (command-R), Step By N (tab), and Run Slower (-) a few times, and it did
> what I expected.
>
>
>>
>> > b) whether the speed is faster, slower, or about the same.
>
>
> How can I measure the speed (which pattern do I run, how many timesteps per
> N, and how big (in pixels) should the main view pane be?)
>
> More importantly, how do I get and build an older version to compare it to?
> Do I need to delete my entire repository or is there a way to tell SVN to
> send me back in time? Which revision should I ask it to give me?

To remove the effects of non-OpenCL code, I tend to do this:
- use a representative pattern like tip-splitting.vti
- set the dimensions to something large, e.g. 512x512x1
- set the timesteps_per_render (running speed) to something large
enough that the number of steps far outweighs the render time, such
that rendering happens once a second or less.

>
> --
>   Robert Munafo  --  mrob.com
>   Follow me at: gplus.to/mrob - fb.com/mrob27 - twitter.com/mrob_27 -
> mrob27.wordpress.com - youtube.com/user/mrob143 - rilybot.blogspot.com
>

--

[Tim Hutton]

The documentation says that clCreateCommandQueue returns
CL_INVALID_VALUE when the properties value isn't supported:
http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clCreateCommandQueue.html
We pass in 0, which should be fine.

Do you get the same error on the CPU on that platform?

--

[Tim Hutton]

On 26 March 2012 23:59, Andrew Trevorrow <and...@trevorrow.com> wrote:
> Tim:
>
>> ...

>> I'm 80% confident that this will fix the problems reported, so if you
>> can test it for me on your machines that would be great. I'm
>> interested in:
>> a) whether everything still works.
>

> I had to add some defines to OpenCL_utils.hpp that are missing from
> Apple's OpenCL headers.  (But I just saw your reply to Robert's report
> so feel free to remove my additions if you think it's better simply
> not to use them in the diagnostics.  I've appended my diagnostic
> output in case it reveals anything interesting.)
>

>> b) whether the speed is faster, slower, or about the same.
>

> A bit slower unfortunately.  Mostly around 10% slower, but pulsate.vti
> dropped from 1800 mcgs to 1200 mcgs.  Still fast enough for me, so
> I'm quite hapopy to live with this if it means Mac users no longer
> see any errors about work group size.

pulsate.vti is slower for me too in the new Ready, but only at smaller
sizes. At e.g. 512x128x32 the new one is faster again.

With this approach the work group size allocation is done by the
OpenCL driver, so we probably will see strange variations. But yes I
think for now we should go with this new approach. In future we can
revisit the issue by looking at performance tuning algorithms (search
the parameter space for the fastest way to process).

>
> Andrew
>
> ------------------------

>
> Found 1 platform(s):
> Platform 1:
> CL_PLATFORM_PROFILE : FULL_PROFILE

> CL_PLATFORM_VERSION : OpenCL 1.0 (Jan  2 2011 18:00:11)

> CL_PLATFORM_NAME : Apple
> CL_PLATFORM_VENDOR : Apple
> CL_PLATFORM_EXTENSIONS :
>
> Found 2 device(s) on this platform.
> Device 1:
> CL_DEVICE_EXTENSIONS : cl_APPLE_gl_sharing

> CL_DEVICE_NAME : ATI Radeon HD 6970M

> CL_DEVICE_PROFILE : FULL_PROFILE
> CL_DEVICE_VENDOR : AMD
> CL_DEVICE_VERSION : OpenCL 1.0
> CL_DRIVER_VERSION : 1.0
> CL_DEVICE_ADDRESS_BITS : 32
> CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE : 0

> CL_DEVICE_MAX_CLOCK_FREQUENCY : 149
> CL_DEVICE_MAX_COMPUTE_UNITS : 10

> CL_DEVICE_NAME : Intel(R) Core(TM) i5-2400 CPU @ 3.10GHz

> CL_DEVICE_PROFILE : FULL_PROFILE
> CL_DEVICE_VENDOR : Intel
> CL_DEVICE_VERSION : OpenCL 1.0
> CL_DRIVER_VERSION : 1.0
> CL_DEVICE_ADDRESS_BITS : 64
> CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE : 64

> CL_DEVICE_MAX_CLOCK_FREQUENCY : 3100
> CL_DEVICE_MAX_COMPUTE_UNITS : 4

> CL_DEVICE_MAX_CONSTANT_ARGS : 8
> CL_DEVICE_MAX_READ_IMAGE_ARGS : 128
> CL_DEVICE_MAX_SAMPLERS : 16
> CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS : 3
> CL_DEVICE_MAX_WRITE_IMAGE_ARGS : 8
> CL_DEVICE_MEM_BASE_ADDR_ALIGN : 1024
> CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE : 128
> CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT : 4
> CL_DEVICE_VENDOR_ID : 16909312
> CL_DEVICE_AVAILABLE : yes
> CL_DEVICE_COMPILER_AVAILABLE : yes
> CL_DEVICE_ENDIAN_LITTLE : yes
> CL_DEVICE_ERROR_CORRECTION_SUPPORT : no
> CL_DEVICE_IMAGE_SUPPORT : yes
> CL_DEVICE_GLOBAL_MEM_CACHE_SIZE : 6291456

> CL_DEVICE_GLOBAL_MEM_SIZE : 12884901888
> CL_DEVICE_LOCAL_MEM_SIZE : 16384
> CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE : 65536
> CL_DEVICE_MAX_MEM_ALLOC_SIZE : 3221225472

> CL_DEVICE_IMAGE2D_MAX_HEIGHT : 8192
> CL_DEVICE_IMAGE2D_MAX_WIDTH : 8192
> CL_DEVICE_IMAGE3D_MAX_DEPTH : 2048
> CL_DEVICE_IMAGE3D_MAX_HEIGHT : 2048
> CL_DEVICE_IMAGE3D_MAX_WIDTH : 2048
> CL_DEVICE_MAX_PARAMETER_SIZE : 4096
> CL_DEVICE_MAX_WORK_GROUP_SIZE : 1
> CL_DEVICE_PROFILING_TIMER_RESOLUTION : 1
> CL_DEVICE_MAX_WORK_ITEM_SIZES : 1, 1, 1
> CL_DEVICE_TYPE : CPU

--

[Andrew Trevorrow]

Robert:

> Failed to open file. Error:
>
> OpenCL_RD::ReloadContextIfNeeded : Failed to create command queue: Invalid value

Someone here

https://discussions.apple.com/thread/3695124?start=0&tstart=0

reported they got the above error when connecting to their Mac remotely
via ssh. Were you doing something like that?

Ready looks like it is calling clCreateCommandQueue() correctly so
I suspect the problem is in your GPU or is due to an Apple-specific bug.

Andrew

[Robert Munafo]

Hi Tim!

On Tue, Mar 27, 2012 at 05:26, Tim Hutton <tim.h...@gmail.com> wrote:

The documentation says that clCreateCommandQueue returns
CL_INVALID_VALUE when the properties value isn't supported:
http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clCreateCommandQueue.html
We pass in 0, which should be fine.

Do you get the same error on the CPU on that platform?

I added printf's to the relevant portion of OpenCL_RD::ReloadContextIfNeeded() to make it look like this:

    // retrieve our chosen device

    {

        const int MAX_DEVICES = 10;

        cl_device_id devices_available[MAX_DEVICES];

        cl_uint num_devices;

        ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ALL, MAX_DEVICES, devices_available, &num_devices);

printf("clGetDeviceIDs(platform_id == %lx, ...) returned num_devices == %d\n",

  ((long) platform_id), ((int) num_devices));

        throwOnError(ret,"OpenCL_RD::ReloadContextIfNeeded : Failed to retrieve device IDs: ");

        if(this->iDevice >= (int)num_devices)

            throw runtime_error("OpenCL_RD::ReloadContextIfNeeded : too few devices available");

        this->device_id = devices_available[this->iDevice];

printf("this->iDevice == %d; setting this->device_id = %lx\n",

  ((int) this->iDevice), ((long) this->device_id));

    }

    // create the context                                          

    this->context = clCreateContext(NULL,1,&this->device_id,NULL,NULL,&ret);

printf("clCreateContext returned this->context == %lx, this->device_id == %lx\n",

  ((long) this->context), ((long) this->device_id));

    throwOnError(ret,"OpenCL_RD::ReloadContextIfNeeded : Failed to create context: ");

    // create the command queue

    this->command_queue = clCreateCommandQueue(this->context, this->device_id, 0, &ret);

    throwOnError(ret,"OpenCL_RD::ReloadContextIfNeeded : Failed to create command queue: ");

Build and Run, then selecting one of the OpenCL patterns, results in the following diagnostic output:

clGetDeviceIDs(platform_id == 7fff0000, ...) returned num_devices == 2

this->iDevice == 0; setting this->device_id = 1021b00

clCreateContext returned this->context == 1016a2600, this->device_id == 1021b00

Seeing that I have 2 devices available, I then added the following in the appropriate place:

 

if (num_devices > 1) {

  this->iDevice = 1;

}

With this change I am now able to run all the patterns on this Sandy Bridge system. Load monitor shows that it is using a single CPU core to perform the calculations.

[Robert Munafo]

No, I'm not ssh'ing to my my laptop. If you read the rest of that discussion thread, you see it's because they're ssh'ing to a "headless server" or a system that has a display card but with no display attached.

I am using my MacBook Pro the way I always do, which is with the MacBook Pro on the desk and me in front of it typing on the keyboard.

> Ready looks like it is calling clCreateCommandQueue() correctly so I suspect the problem is in your GPU or is due to an Apple-specific bug.

Yes, that's probably essentially right. I think it's failing because we're trying to do something on the Sandy Bridge GPU that isn't supported by the Apple's OpenCL drivers or libraries. See my later message, where I insert printf's, discover that there are two devices, and "solve the problem" by using the CPU as the OpenCL compute device.

There are going to be a lot of users on Sandy Bridge and Ivy Bridge systems who have no discrete GPU, so we really ought to work on the Intel integrated GPU.

On Tue, Mar 27, 2012 at 07:20, Andrew Trevorrow <and...@trevorrow.com> wrote:

Robert:

> Failed to open file. Error:
>
> OpenCL_RD::ReloadContextIfNeeded : Failed to create command queue: Invalid value

Someone here

https://discussions.apple.com/thread/3695124?start=0&tstart=0

reported they got the above error when connecting to their Mac remotely via ssh.  Were you doing something like that?

Ready looks like it is calling clCreateCommandQueue() correctly so I suspect the problem is in your GPU or is due to an Apple-specific bug.

Andrew

[Tim Hutton]

There's a function on the Action menu: Select OpenCL Device... that
should show the two devices and allow you to select which one to use.
Did you try that?

It's very odd that the 8-threads i7-2720QM should report
CL_DEVICE_MAX_WORK_GROUP_SIZE of 1. That's why it's only using one
core. Is there any way to update your OpenCL drivers?

--

[Robert Munafo]

Hi Tim,

On Tue, Mar 27, 2012 at 11:14, Tim Hutton <tim.h...@gmail.com> wrote:

There's a function on the Action menu: Select OpenCL Device... that
should show the two devices and allow you to select which one to use.
Did you try that?

 

Oh, that's nice. I didn't see that command. That means I can test the other card on my Mac Pro. Very good idea.

All right, I just took out the debugging hacks and built again. When I invoke the "Select OpenCL Device" command, I see:

  Apple : ATI Radeon HD 6750 M

  Apple : Intel(R) Core(TM) i7-2720QM CPU @ 2.20 GHz

with the first item selected by default. If I select the second item, then yes, it works, and it's about 5 times faster because it's using all of the cores.

If I select the first item again, I get the errors.

I do indeed have an ATI Radeon HD 6750 M GPU in this computer, but I keep it turned off to save battery life, using the excellent and very popular gfxCardStatus utility.

Apple's system software is very aggressive about using the discrete GPU even when it is absolute overkill, and even after all client apps have been shut down.

Even something as trivial as the Twitter dashboard widget (really!) causes the OS to turn on the Radeon GPU. The battery life is about 2.5 times longer if you disable theis behaviour. See:

  http://www.anandtech.com/show/4205/the-macbook-pro-review-13-and-15-inch-2011-brings-sandy-bridge/9

  http://reviews.cnet.com/8301-13727_7-20042341-263.html

  http://www.thepunditreport.com/2011/08/revealed-mac-os-x-lion-battery-drain.html

  http://tidbits.com/article/11982

  http://www.macworld.com/article/1163166/gfxcardstatus_2_1.html

  http://mihail.stoynov.com/2011/08/18/dramatically-increase-battery-life-on-a-macbook-pro-with-2-gpus/

In the rare event that I want to play a 3D game I then turn it on manually using the gfx menubar icon.

- Robert

It's very odd that the 8-threads i7-2720QM should report
CL_DEVICE_MAX_WORK_GROUP_SIZE of 1. That's why it's only using one
core. Is there any way to update your OpenCL drivers?

[Tim Hutton]

On 27 March 2012 17:09, Robert Munafo <mro...@gmail.com> wrote:
> Hi Tim,
>
> On Tue, Mar 27, 2012 at 11:14, Tim Hutton <tim.h...@gmail.com> wrote:
>
>> There's a function on the Action menu: Select OpenCL Device... that
>> should show the two devices and allow you to select which one to use.
>> Did you try that?
>
>
> Oh, that's nice. I didn't see that command. That means I can test the other
> card on my Mac Pro. Very good idea.
>
> All right, I just took out the debugging hacks and built again. When I
> invoke the "Select OpenCL Device" command, I see:
>
>   Apple : ATI Radeon HD 6750 M
>   Apple : Intel(R) Core(TM) i7-2720QM CPU @ 2.20 GHz
>
> with the first item selected by default. If I select the second item, then
> yes, it works, and it's about 5 times faster because it's using all of the
> cores.
>
> If I select the first item again, I get the errors.
>
> I do indeed have an ATI Radeon HD 6750 M GPU in this computer, but I keep it
> turned off to save battery life, using the excellent and very popular
> gfxCardStatus utility.

Have you tried turning it on and then running Ready?

>
> Apple's system software is very aggressive about using the discrete GPU even
> when it is absolute overkill, and even after all client apps have been shut
> down.

What is a 'discrete GPU'?

>
> Even something as trivial as the Twitter dashboard widget (really!) causes
> the OS to turn on the Radeon GPU. The battery life is about 2.5 times longer
> if you disable theis behaviour. See:
>
>   http://www.anandtech.com/show/4205/the-macbook-pro-review-13-and-15-inch-2011-brings-sandy-bridge/9
>   http://reviews.cnet.com/8301-13727_7-20042341-263.html
>   http://www.thepunditreport.com/2011/08/revealed-mac-os-x-lion-battery-drain.html
>   http://tidbits.com/article/11982
>   http://www.macworld.com/article/1163166/gfxcardstatus_2_1.html
>   http://mihail.stoynov.com/2011/08/18/dramatically-increase-battery-life-on-a-macbook-pro-with-2-gpus/
>
> In the rare event that I want to play a 3D game I then turn it on manually
> using the gfx menubar icon.
>
> - Robert
>
>
>> It's very odd that the 8-threads i7-2720QM should report
>> CL_DEVICE_MAX_WORK_GROUP_SIZE of 1. That's why it's only using one
>> core. Is there any way to update your OpenCL drivers?
>
>
> --
>   Robert Munafo  --  mrob.com
>   Follow me at: gplus.to/mrob - fb.com/mrob27 - twitter.com/mrob_27 -
> mrob27.wordpress.com - youtube.com/user/mrob143 - rilybot.blogspot.com
>

--

[Robert Munafo]

> I do indeed have an ATI Radeon HD 6750 M GPU in this computer, but I keep it turned off to save battery life, using the excellent and very popular gfxCardStatus utility. 

 

Have you tried turning it on and then running Ready?

I wouldn't do that. If it requires having half the battery life, or less, then I'll just not use the application.

Other applications, like the Flash plug-in in the Chrome browser or the Twitter widget I mentioned earlier, also try to use the discrete GPU, but they successfully detect that it is unavailable and connect to the integrated GPU instead. (I don't know how they detect this or what they're actually using a GPU for, but it's probably not OpenCL.)

> Apple's system software is very aggressive about using the discrete GPU even when it is absolute overkill, and even after all client apps have been shut down.

What is a 'discrete GPU'?

A discrete GPU, in this context, means a GPU that is in its own package (silicon die and packaging attached to the motherboard) rather than being "integrated" into the same package as the CPU (as with Intel CPUs aimed at the mobile market, beginning with the "Clarkdale" Core i3 and i5 in 2009)

 >> It's very odd that the 8-threads i7-2720QM should report CL_DEVICE_MAX_WORK_GROUP_SIZE of 1. That's why it's only using one core. Is there any way to update your OpenCL drivers?

I forgot to answer this one. It must have been using only 1 thread simply because I hacked the code to set this->iDevice to 1 at a point in time when the program had (presumably) already chosen its settings based on device 0. It was just a temporary test to see if I could get device 1 to work, because I didn't know about that menu command.

[Andrew Trevorrow]

Robert:

> I do indeed have an ATI Radeon HD 6750 M GPU in this computer, but I keep
> it turned off to save battery life, using the excellent and very
> popular gfxCardStatus utility.

Tim:

> Have you tried turning it on and then running Ready?

Robert:

> I wouldn't do that. If it requires having half the battery life, or less,
> then I'll just not use the application.

I think Tim just wanted to know if turning on the GPU prevents the
error messages Ready was reporting (when the GPU is selected as your
OpenCL device of course). Once you've done that test you can select
the CPU as your OpenCL device and Ready will remember that setting
in your ReadyPrefs file.

Tim, I'm not sure if this problem is common enough to worry about,
but maybe the code in OpenCL_RD::ReloadContextIfNeeded() should
automatically switch to another device (if there is one, and with
a suitable warning) if either clCreateContext or clCreateCommandQueue
return an error.

Andrew

[Robert Munafo]

Oh yes, sorry. It works on the ATI Radeon GPU if I have it turned on.

I think instead of trying to automatically handle this error, (which seems like us trying a bit too hard to second-guess things, and might even violate some kind of OpenGL programming style guideline), perhaps we can just give the user a better error message. The thing we print now is nearly useless. I suggest an error message more like the following:

I cannot use "Apple : ATI Radeon HD 6750 M GPU" to open the file "schlogl.vtl". You might wish to try selecting a different OpenCL Device or to open a different pattern file.

Details of error:

OpenCL_RD::ReloadContextIfNeeded : Failed to create command queue: Invalid value

See? Much more informative, without losing the technical details.

- Robert

On Tue, Mar 27, 2012 at 21:10, Andrew Trevorrow <and...@trevorrow.com> wrote:

I think Tim just wanted to know if turning on the GPU prevents the
error messages Ready was reporting (when the GPU is selected as your
OpenCL device of course).  Once you've done that test you can select
the CPU as your OpenCL device and Ready will remember that setting
in your ReadyPrefs file.

Tim, I'm not sure if this problem is common enough to worry about,
but maybe the code in OpenCL_RD::ReloadContextIfNeeded() should
automatically switch to another device (if there is one, and with
a suitable warning) if either clCreateContext or clCreateCommandQueue
return an error.

[Tim Hutton]

We've had a go at fixing all the errors reported for 0.2, so let's cut
a bugfix release 0.2.2 and make new distributables for testing. I've
changed the version number and added a new section in changes.html. Is
there anything else to go in?

[Andrew Trevorrow]

Tim:

Nope, looks good to me.

Andrew