OpenCL support for Mali-T628 MP6 on Arndale Octa?

Thanks for the reply Chris,

We are using Exynos 5420 board from Insignal. Yes it has kernel level binary inside the device. Getting a driver at Android usespace will be very helpful, because we run our application on different vendor SOCs to compare the performance. The other SOCs we have has Android as the OS.

I will check in Insignal forums.

Thanks for your help.

Veeranna

Thanks for the reply Chris,

We are using Exynos 5420 board from Insignal. Yes it has kernel level binary inside the device. Getting a driver at Android usespace will be very helpful, because we run our application on different vendor SOCs to compare the performance. The other SOCs we have has Android as the OS.

I will check in Insignal forums.

Thanks for your help.

Veeranna

Hi Chris,

I got the driver from Samsung/Insignal, and able to build the Andorid Image and Run. But we are observing random behavior in OpenCL APIs. Some time clcreateKernel fails and if re-run the app clcreateKernel returns success but clcreatebuffer fails with error code as zero and returned memory pointer as NULL. 

Can you give some to hint where to look.

Thanks,

Veeranna

Hi Veeranna,

Can you provide us with a simple reproducer? Also please let me know the driver version you're using, with adb pull /vendor/lib/egl/libGLES_mali.so && strings libGLES_mali.so | grep r[0-9]p[0-9] assuming it's in vendor and not system.

Thanks,

Chris

Hi Chris,

With above said command, I got version as 1.4 Midgard-"r3p0-01bet0. I will try to give simple app to run.

Thank you for help.

Veeranna

Interestingly simple median filter example runs fine(openCL on T628), but our application fails. Our application has many kernels and huge memory bandwidth will it be the reason?

Any suggestion for debugging will be helpful.

Thanks,

Veeranna

Hi veerannah,

I doubt simply having a lot of kernels would in itself be a problem, nor would I expect memory bandwidth to be an issue. Can you confirm whether you are creating all of your kernels up-front before any clEnqueueNDRange commands take place, or do you somewhat interleave kernel compilation and execution? Can you also confirm whether all program objects compile without error prior to the calls to clCreateKernel that consume them?

Thanks,

Chris

Hi Chris,

We create all our kernels first, then we allocate required buffers and then call clEnqueueNDRange. Sometime it fails to create some kernels(some are got created). And sometimes it fails in buffer allocation. It didnt hit clEnqueueNDRange yet.

Thanks,

Veeranna

It's interesting that its not deterministic where it fails. Without a reproducer I'm afraid I'm just guessing at possible causes. Also the driver itself is quite old now, we're currently at r4p0, so it is worth asking them when they plan to provide an up to date version as this could fix it.

Thanks,

Chris

Hi Chris,

Finally we are able to run our application on T628. But performance numbers are not good. Do we get any improvement if we move r4p0 driver?

Any other suggestions to improve the GPU performance will be helpful.

Thanks,

Veeranna

Hi Veeranna,

There are a number of performance improvements present in the r4p0 driver not present in previous releases. Keep in mind that OpenCL is not performance portable, so an application optimized for another platform, or otherwise written with another architecture in mind, may not be performant when run on another platform/architecture. The below materials contain advice and detail some of the differences and considerations when moving from desktop to Mali, so let us know if they helps or if you have any further queries and we'll be happy to help.

There is the Developer Guide: Mali-T600 Series GPU OpenCL Developer Guide « Mali Developer Center

There is also the OpenCL faq: http://malideveloper.arm.com/downloads/OpenCL_FAQ.pdf

And the Laplace case study by timhar01 Technical presentation about ARM Mali-T600 GPU and ARM Mali-T700 GPU Compute - YouTube (although I recommend you watch the whole video)

Hope this helps,

Chris

Hi Chris,

Where can we find benchmark numbers for examples kernels provided with SDK?. We ran "sobel" filter example for 1024x1024 image, profile number measured across kernel trigger and wait measures as 5msec. And Laplace transform for 4096x2048 image, number measures as 10msec. Is these numbers expected for for Mali T628 @ 533Mhz GPU? We measured the numbered across clEnqueueNDRange() and immediate clFinish().

Thanks,

Veeranna

Hi Veerannah,

We don't have stock performance numbers for the SDK samples to hand.  However what I will do later this week is run Sobel and Laplace on a platform here and get some numbers back to you.

HTH, Tim

Hi Tim,

We really appreciate your quick reply. Looking forward to see the performance numbers.

Thanks,

Veeranna

Hi Veeranna,

I ran the Sobel filter as provided by the Mali OpenCL SDK, modifying it to run a 1024x1024 image and I see the following...

Queued time:    0.11ms

Wait time:        0.222292ms

Run time:         0.680683ms

The Mali OpenCL SDK doesn't include a Laplace transform... which one are you using?

Best regards, Tim

Hi Veeranna,

I realise I didn't include the platform details.  I'm running on a Mali-T604 platform with the GPU clocked @ 533MHz.  I'm using a later version of the drivers than you, I believe, and that would likely mean the performance you would expect on your platform is slightly worse.  On balance though I would expect the performance on your Mali-T628 system to be roughly equivalent to mine.

HTH, Tim

Hi Tim,

Thank you for providing information on performance number. With "cl_event" profiling we are getting numbers as

Queued time:    0.11ms

Wait time:        0.2725ms

Run time:         0.7586ms

But if we measure numbers across clEnqueueNDRangeKernel() and clfinish() using gettimeofday we are seeing number is 1.5msec. Is this API/driver overhead?

Thanks,

Veeranna