Archives Discussions

[gigo]

The last few days I was chasing a bug which caused all kinds of strange and random driver crashes..

I finally tracked it down to me giving invalid arguments to clEnqueueNDRangeKernel(), namely the global and local work size.

In some cases I accidentally passed a global work size which was not evenly divisable by the local work size. Of course, that is an error on my end, but according to the standard the function should return an error if something like that is attempted. As far as I can see it never actually does that for the not-evenly-divisable error.

The documentation says:

CL_INVALID_WORK_GROUP_SIZE if local_work_size is specified and number of work-items specified by global_work_size is not evenly divisable by size of work-group given by local_work_size [...].

With all the different limitations on the global and local work size I feel like this must be one the easiest errors to check for. So I wonder why the OpenCL implementation happily accepts it and returns CL_SUCCESS. Can someone clarify whether this is actually a bug in the OpenCL implementation?

My system specs:

Windows 8.1 x64

Catalyst 15.7 (OpenCL 2.0 AMD-APP (1800.5)

Radeon HD 7870

Intel Core i5-2500K

[vladimir_1 · Accepted solution]

OpenCL 2.0 allows non-uniform workgroup sizes.

http://developer.amd.com/community/blog/2014/12/11/opencl-2-0-generic-address-space-program-scope-va...

so the CL_SUCCESS is an expected behaviour.

Driver crash is not ofcause 😃

[gigo]

Oh thanks, I didn't know that.

The version on the platform reports "OpenCL 2.0 AMD-APP (1800.5)" but on the device it is just "1.2".

According to Wikipedia my GPU (Pitcairne) should support OpenCL 2.0, is that correct?

The really bad thing is that the driver does not crash immediatly but randomly after some of these non-evenly divisable work sizes have been enqueued. Some of these kernels actually do get invoked and do something but the results we're not correct. I would guess that the non-uniform workgroup at the end did not run (correctly).

[vladimir_1]

Documentation for clEnqueueNDRangeKernel says that you should pass –cl-uniform-work-group-size option when compiling the kernel to raise  CL_INVALID_WORK_GROUP_SIZE error

[gigo]

Ok, again something I didn't notice, thanks. I also realized that I did not always check which version of the OpenCL documentation I was looking at, e.g. the excerpt in my first post was from version 1.2. Now that I think about it, what version should I refer to since my platform is 2.0 and the device reports 1.2?

But the behaviour I got is buggy nonetheless. Either non-uniform work sizes are supported and executed correctly or not, not something in between. In my case the driver accepted the non-uniform work size and partially executed it but crashed some time later.

[dipak]

Hi,

When I checked a similar issue as reported here Kernel Workspaces , I didn't observe any crash for "not divisible work-groups". Could you please provide a reproducible test-case?

Regards,

[gigo]

Hi,

I must admit that the error was on my end, all-clear on the driver side 😉

The driver crash was caused by repeated out-of-bounds access to a buffer, because I made an error during the index calculation.

My observation that the execution of the non-uniform workgroup at the end has problems was also incorrect. My kernel uses local memory which is not initialized correctly on the non-uniform workgroup. The work items that do run in this group then work on uninitialized memory which is obviously not a good idea. This caused the results to deviate from what I expected and I blamed it on the driver, oops ..

Thanks for the great support anyway!