OpenCL on Sony Xperia™ devices

Boost the performance of your Android app with OpenCL

Although our Xperia™ devices have become much more powerful the last couple of years, there might still be cases when your application needs more processing power. With OpenCL™, you can use the power of the GPU to handle resource intensive tasks in your app. This article is a short introduction to OpenCL, and how to get it up and running on your Sony Xperia device.

Jim Rasmusson
Jim Rasmusson, Master Researcher at Sony Mobile.

In this article, Jim Rasmusson, Master Researcher, will take you through the steps you need to take to get started using OpenCL* in your app. A code example is also provided, see OpenCL code example page. But first, we’ll explain some basics about OpenCL.

What is OpenCL?

The graphics processors (also called GPUs) inside the Xperia devices have taken a big leap forward this year with a new hardware architecture that enables so called General Purpose Graphics Processing Unit (GPGPU) capabilities. In essence, this means that the GPU can be used as a powerful array processor for general purpose processing. In our current top models, the Xperia Z1 and Xperia Z Ultra, the GPU actually has 128 small processing cores in total.

In order to program these array processors efficiently, a new programming framework is needed. One option is to use OpenCL, which has been around a couple of years for desktop computers. OpenCL has been created by Khronos™ as an open standard for high performance computations. It is now quite mature, has good industry support and an interesting roadmap.

OpenCL can be particularly useful when running data parallel processing tasks. Among the most efficient workloads are various types of multimedia processing tasks like image, camera and computer vision type processing. But you may also find benefits with other types of data parallel tasks like encryption, or when simulating physics or particle systems.

OpenCL is available in several of our Sony Xperia devices, including the Xperia Z, Xperia ZL, Xperia Tablet Z and Xperia ZR (in addition to Xperia Z1 and Xperia Z Ultra mentioned above).

The OpenCL platform model
OpenCL Platform Model
The OpenCL platform model: You have one Host and one or more Compute Devices. A Compute Device has one or more Compute Units, which have one or more Processing Elements. (Image supplied courtesy of Khronos Group)

We currently support OpenCL version 1.1**, and our OpenCL implementation adheres to the OpenCL Platform Model (see the illustration above). We have a Host, running on the CPU, and a Compute Device, that, in our implementation, is the GPU. For example, in Xperia Z1 and Xperia Z Ultra, the GPU is the Adreno 330, which has 4 Compute Units, each with 32 Processing Elements, giving us 128 processing elements in total.

OpenCL consists of different parts: the platform and runtime APIs, and the actual kernel programming language itself, OpenCL-C, which is based on C99. If your favourite C code is not overly complicated, it is often fairly straightforward to port it to OpenCL-C. You use the platform and runtime APIs to query, setup and start execution on the OpenCL enabled subsystems available in your system (for example, the GPU, and maybe even the CPU).

The OpenCL code example
To be able to use this example, you should be familiar with C, C++, and Java programming in an Android environment.

In our small image noise reduction code example, we will use a bilateral filter to filter an image. This is quite math heavy. As we will see, the OpenCL version of this code executes around 60 times faster (on some of our Xperia devices) in OpenCL compared to the (single threaded) reference C code!

Since OpenCL is not natively supported in Android, we have to use the Android NDK. If you take a look in the code example project, you will see the typical files necessary for an NDK project, such as the Android.mk file, and the jni folder with your C and C++ files. In the asset folder, there is also a bilateralKernel.cl file, a text file that contains our OpenCL-C kernel code.

If we take a look at the OpenCL–C kernel code below, we see that it is quite “C like”, but with some additional data types and a few extra commands. In particular, note the vector data types (for example float4), and their associated nice vector component selection syntax. In the example below we use currentPixel.xyz to select the first three channels (rgb) of this particular float4 variable.

kernel void bilateralFilterKernel(__global uchar4 *srcBuffer, __global uchar4 *dstBuffer, const int width, const int height){
     int x = get_global_id(0);
     int y = get_global_id(1);
     int centerIndex = y * width + x;
     float4 sum4 = (float4)0.0f;

          if ( (x >= filterWidth) && (x < (width - filterWidth)) && //avoid reading outside of buffer
              (y >= filterWidth) && (y < (height - filterWidth)) )
        {
               float4 centerPixel = oneover255 * convert_float4(srcBuffer[centerIndex]);
               float normalizeCoeff = 0.0f;

              for (int yy=-filterWidth; yy<=filterWidth; ++yy)
             {
                       for (int xx=-filterWidth; xx<=filterWidth; ++xx)
                      {
                                 int thisIndex = (y + yy) * width + (x + xx);
                                float4 currentPixel = oneover255 *convert_float4(srcBuffer[thisIndex]);
                                float domainDistance = fast_distance((float)(xx), (float)(yy));
                               float domainWeight = exp(-0.5f * pow((domainDistance/sigmaDomain),2.0f));

                               float rangeDistance = fast_distance(currentPixel.xyz, centerPixel.xyz);
                              float rangeWeight = exp(-0.5f * pow((rangeDistance/sigmaRange),2.0f));

                              float totalWeight = domainWeight * rangeWeight ;
                             normalizeCoeff += totalWeight;

                             sum4 += totalWeight * currentPixel;
                      }
              }
              sum4 /= normalizeCoeff;
      }
      dstBuffer[centerIndex] = convert_uchar4_sat_rte(255.0f * sum4);
}

One example function built-into OpenCL is distance. It calculates the Euclidean distance between two variables and is handy for image processing, for example.

So

float rangeDistance = distance(currentPixel.xyz, centerPixel.xyz);

corresponds to the C code

float rangeDistance = 0.0f;
for (int c=0; c<3; ++c)
                  rangeDistance += (currentPixel[[c]] - centerPixel[[c]]) * (currentPixel[[c]] - centerPixel[[c]]);

In our OpenCL kernel code example, we have used the faster version of distance“fast_distance” , which brings even faster execution speeds. To get the typical bilateral filter effect (and associated image noise reduction), it is necessary to iterate this filter a few times over the image. This why we call the same bilateral filter routine three times in this example.

NativeC bilateral filter
Left: noisy original image. Right: after bilateral filter with NativeC
OpenCL bilateral filter
Left: noisy original image. Right: after bilateral filter with OpenCL. Note that the processing time is now 51 ms compared to 3022 ms in the NativeC example.

If you have an Android device that supports OpenCL, you can run the supplied openclexample.apk file to check how it works. You can also see that the execution speed is much faster using OpenCL on the GPU when compared to the plain single threaded c-code running on the CPU (tested on Sony Xperia Z1). In addition to the speed benefit, you may also find that you decrease energy consumption by utilizing OpenCL on the GPU compared to using standard programming methods on the CPU.

OpenCL actively demonstrates that the GPU units in recent mobile devices are more than capable of handling non-graphical computing across a variety of apps. Keep OpenCL in mind when considering how to enhance flexibility, functionality, and performance of your apps, and let us know in the comments below if you have any further questions!

More information

• Read more about OpenCL
• Check out Android Developer to learn more about programming for Android
• Get the Android NDK

*OpenCL and the OpenCL logo are trademarks of Apple Inc. used by permission by Khronos.
**Product is based on a published Khronos Specification, and is expected to pass the Khronos Conformance Testing Process. Current conformance status can be found at www.khronos.org/conformance.

Comments 26

Sort by: Newest | Oldest | Most popular

  1. By Jory Ferrell

    #1

    How would I go about creating and passing an efficient 3d buffer for non-image based data?

    • By Jory Ferrell

      #2

      For instance, how would I create a 256^3 buffer of floats?

  2. By Olivier Valery

    #3

    Hi

    My question is related to the sony Xperia Z5. As said in the previous discution, it is compatible with OpenCL 1.2. However the example provided on this web site only concern opencl 1.1. Do you know wher to find the libopencl.so that match with the sony xperia Z5 hardware ?

    Best regards,

    Olivier

  3. By Olivier Valery

    #4

    I have a question:

    I use OpenCL on Sony xperia Z Ultra for 2 years already (for a research purpose). I would to “update” my hardware and buy a new sony xperia device. Do the new xperia Z5 support opencl ? In term of GPU computing performance, which one would you advice me ?

    thanks,

    Best regards,

    Olivier

    • By Jim Rasmusson

      #5

      Hi Olivier

      Yes, the Z5 models are currently the best for OpenCL GPU processing.

      If you are asking specifically which of the three Z5 model is best (compact, standard or premium), I am not sure. I would guess they are fairly similar. If this is important for you, let me know and I can find out.

      Thanks and Best Regards!
      /Jim

      • By Olivier Valery

        #6

        Thank you for your answer. Can you provide me their hardware characteristics in term of max group size (opencl), local memory …

        thanks a lot,
        olivier

        • By Jim Rasmusson

          #7

          Hi again Olivier

          Sorry for late response! (just saw your follow up question)

          The Z5s are build using the Snapdragon 810 chipset from Qualcomm. The GPU, Adreno 430, comes with full profile OpenCL 1.2 support (which is great news if you are moving your desktop CL projects to your phone).
          The local RAM is increased to 32kB (compared to 8KB in Adreno 3xx) per compute unit. Max workgroup size is 1024. In general Adreno 430 has improvements all over the place compared to 330.

          Hope this helps.

          Thanks!
          /Jim

  4. By Rob V

    #8

    Hi Jim,

    i want to run my own openCL program on a Xperia Z3. I’ve a little problem with structs. Does the Z3 support structs as argument to a kernel?

    The problem ist, when i pass an array of structs to the kernel, just the first element ist set right in the kernel.

    do you know about this?

    thanks a lot
    robin

    • By Jim Rasmusson

      #9

      Hi Robin

      Yes you can use structs in OpenCL on the Xperia phones. But array of structs can be tricky (especially with alignment).

      In general there are some restrictions (described in section 6.8 of the OpenCl 1.1 specification).

      It is adviced to use “cl_” when declaring variables on the host side (to make sure the same amount of data is allocated on both host and device). Also having the largest types first in your struct may help and enlargening your types to consume 32 or 128 bits may help getting the alignment correct (e.g., use float4 instead of float3)

      There is a code example in the Adreno SDK, “ParticleSystemCLGLES”, that make use of structs. Not sure if it address your particular problem but check it out.

      Good luck!
      /Jim

  5. By Jonathan Reese

    #10

    Hi,

    Thank you for the great tutorial. It has been very helpful to start doing some OpenCL on Android!

    My issue is OpenGL/OpenCL interop. I currently have a GL texture, read it into a buffer with glReadPixels, handle it through OpenCL (using your example) and then send it back to GL for rendering. This is of course less than ideal, but it works.

    I was not able to make OpenGL/CL interop work (yet) on Android. I have been searching on such examples but only desktop examples are present. One of big hurdles was the EGL context. I think I fixed that. However, whenever I try to create a CL texture from GL one (with clCreateFromGLTexture2D) I get a CL_INVALID_IMAGE_FORMAT_DESCRIPTOR.

    Any pointers or examples of working code on Android would be awesome.

    Regards.

    • By Jim Rasmusson

      #11

      Hi Jonathan,

      I am glad you found the tutorial useful!

      CL/GL interop can be a bit tricky on Android when EGL is involved. Check Section 9.8 in the CL1.1 spec, which contains info on image formats and additionally describes some potentially helpful query functions for debugging (such as clGetGLObjectInfo and clGetGLTextureInfo). Also Section 9.7 provide some potentially helpful info.

      In general, since OpenCL is not natively supported on Android, there is a scarse supply of tutorials and open source projects. However there is one; Amir Chohan has been kind enough to open source his project using OpenCL on Android. https://github.com/amirchohan/HDR

      Additionally, the Adreno SDK contains some OpenCL documents and samples.

      Good luck!

      /Jim

  6. By Olivier Valery

    #12

    Thank you for your article that open new opportunity for programmers.

    I try many times to run a barrier on my Sony Android device without any success. Do this function supported ?

    • By Jim Rasmusson

      #13

      Hi Olivier,

      Yes barriers are supported. Can you provide a code example where it fails?

      Thanks
      Jim

  7. By A. Bruenner

    #16

    Assuming a stock Android 4.4 (on Xperia Z1, not rooted):

    1. Is there also support for OpenCL 1.2 (despite the OpenCL version used in the tutorial, which was 1.1)?

    2. Is the libOpenCL.so on the device currently of version 1.2?

    3. How would I retrieve the libOpenCL.so from the device?
    (I tried pulling the lib file from the device with adb, but of course do not have root access to the directories in question)

    4. If the libOpenCL on the device is not of version 1.2, where could I get that lib?

    Any help greatly appreciated.

    • By Jim Rasmusson

      #17

      Hi A Bruenner

      The chipset in the Xperia Z1 supports OpenCL 1.1 only. It will not be updated to OpenCL1.2 as the hardware does not fulfil the requirements.

      libOpenCL.so can be found in vendor/lib/ on the phone (use adb shell to inspect the file system on the phone). You can “pull” (copy) the libOpenCL.so from the phone to your PC using “abd pull” (also on a non-rooted phone). Note that the OpenCL driver stack is made specifically for this chipset/phone (hardware plus software). It most likely does not work if you move or replace these libs to/from other phones.

      Hope this answers your questions?

      /Jim Rasmusson

      • By A. Bruenner

        #18

        Thanks for your reply and your work on enabling/supporting OpenCL on Sony mobile devices!

    • By A. Bruenner

      #19

      Point 1-3: Resolved.

  8. By Eugenio Culurciello

    #20

    Nexus 5 does not have libopencl.so in /system/lib
    I guess we cannot do anything without that, right?

    Or did you link the provided libopencl.so into the project?
    Cannot be found for me 😉

    • By Jim Rasmusson

      #21

      Hi Eugenio,

      Unfortunately I think OpenCL is not supported in Google’s Nexus devices. You can check also in /vendor/lib/ for libOpenCL.so, but if it is not there you are probably out of luck.

      /Jim

  9. By Eugenio Culurciello

    #22

    Thank you Sergej, this is great. I am trying this on a Nexus 5 and S4. It looks like I have to re-compile the opnCL lib for these devices right?
    Your example does not work for me, it reports runtime errors:
    12-17 16:35:31.841: E/AndroidRuntime(27712): FATAL EXCEPTION: main
    12-17 16:35:31.841: E/AndroidRuntime(27712): Process: com.sony.openclexample1, PID: 27712
    12-17 16:35:31.841: E/AndroidRuntime(27712): java.lang.IllegalStateException: Could not execute method of the activity
    12-17 16:35:31.841: E/AndroidRuntime(27712): at android.view.View$1.onClick(View.java:3814)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at android.view.View.performClick(View.java:4424)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at android.view.View$PerformClick.run(View.java:18383)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at android.os.Handler.handleCallback(Handler.java:733)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at android.os.Handler.dispatchMessage(Handler.java:95)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at android.os.Looper.loop(Looper.java:137)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at android.app.ActivityThread.main(ActivityThread.java:4998)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at java.lang.reflect.Method.invokeNative(Native Method)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at java.lang.reflect.Method.invoke(Method.java:515)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at com.android.internal.os.ZygoteInit$MethodAndArgsCaller.run(ZygoteInit.java:777)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at com.android.internal.os.ZygoteInit.main(ZygoteInit.java:593)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at dalvik.system.NativeStart.main(Native Method)
    12-17 16:35:31.841: E/AndroidRuntime(27712): Caused by: java.lang.reflect.InvocationTargetException
    12-17 16:35:31.841: E/AndroidRuntime(27712): at java.lang.reflect.Method.invokeNative(Native Method)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at java.lang.reflect.Method.invoke(Method.java:515)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at android.view.View$1.onClick(View.java:3809)
    12-17 16:35:31.841: E/AndroidRuntime(27712): … 11 more
    12-17 16:35:31.841: E/AndroidRuntime(27712): Caused by: java.lang.UnsatisfiedLinkError: Native method not found: com.sony.openclexample1.OpenCLActivity.runNativeC:(Landroid/graphics/Bitmap;Landroid/graphics/Bitmap;[I)I
    12-17 16:35:31.841: E/AndroidRuntime(27712): at com.sony.openclexample1.OpenCLActivity.runNativeC(Native Method)
    12-17 16:35:31.841: E/AndroidRuntime(27712): at com.sony.openclexample1.OpenCLActivity.showNativeCImage(OpenCLActivity.java:132)
    12-17 16:35:31.841: E/AndroidRuntime(27712): … 14 more

  10. By Anonymous

    #23

    Hi,
    Nice post. Is the Xperia Z1 comes with OpenCL support out of the box or there is a need
    to install something/drivers so that it would run?

    thanks

    • By Anonymous

      #24

      Hi

      Yes, it comes with the OpenCL drivers built-in. But you need libOpenCL.so somewhere in your build environment on your pc to be able to build for the Xperia devices. See my code example.

      /Jim

      • By zhengjun he

        #25

        Hi Jim,
        I download your example code,but when build the native code with the tool NDK in windows,a error is prompted as below:
        D:\openCL\openclexample1\jni>ndk-build
        Android NDK: WARNING:D:/openCL/openclexample1//jni/Android.mk:openclexample1: non-system libraries in linker flags: D:/openCL/openclexample1//jni/../extra_libs/libOpenC
        Android NDK: This is likely to result in incorrect builds. Try using LOCAL_STATIC_LIBRARIES
        Android NDK: or LOCAL_SHARED_LIBRARIES instead to list the library dependencies of the
        Android NDK: current module
        [armeabi-v7a] Install : libopenclexample1.so => libs/armeabi-v7a/libopenclexample1.so

        When I build the apk and run it on android phone,even the native C function will crash once it starts to process due to the same reason as posted by Eugenio’s first post

        Coud you pls help me how to solve this problem?

        • By Jim Rasmusson

          #26

          Hi Zhengjun

          Sorry for late response (it has been kinda low activity on the comment front so I have not checked here as of late).

          Despite the warnings libopenclexample1.so seems to build OK.. ([armeabi-v7a] Install : libopenclexample1.so => libs/armeabi-v7a/libopenclexample1.so)

          What phone do you have? and firmware version?

          /Jim

1-26 of 26 comments.