<p>Mobile devices are carving their niche into the world of computing with more processing power day by day. GPUs on mobile devices have been around for a while, but using them for accelerating computation is still quite new. Until recently, the only way to access the GPU was through OpenGL. Around december 2008, Khronos released OpenCL, a generic API for accelerating non-graphics tasks. OpenCL enables us to take advantage of acceleration hardware. Since it is an open standard, many hardware vendors provide support on their devices. With the recent release of Adreno and Mali SDKs, you can now run OpenCL code on mobile GPUs.</p>
<p>Today's post is going to be about how to do image processing on camera feed on Android devices using OpenCL. To successfully complete this tutorial, you need to have basic knowledge of how to write OpenCL kernels and native code for Android.</p>
<h3 id="opencl">OpenCL</h3>
<p>To use OpenCL on your mobile device's GPU, you must obtain the compatible libOpenCL.so library and the corresponding OpenCL headers.</p>
<p>I initially thought I could just use the C++ OpenCL header available from Khronos website for development, however that didn't work. I faced compilation errors because of incompatible OpenGL headers. There are couple of ways to get the compatible headers. You can either use the <a href="https://github.com/arrayfire/androidcl/tree/master/jni/include">headers</a> provided along with the example application or extract them from Adreno/Mali SDK. For Adreno SDK, you can find the headers at <code>&lt;Adreno_SDK&gt;/Development/Inc/CL</code>. For Mali SDK, you can find the headers at <code>&lt;MALI_SDK&gt;/include/CL</code>.</p>
<p>For newer devices such as Samsung Galaxy Note 3 and Samsung Galaxy S5, it is very likely the library is already present on the mobile device. My test device is a Note 3 for which i found the library under <code>/system/vendor/lib</code>. You can do <code>adb pull /system/vendor/lib/libOpenCL.so ./</code> to get the file to your desktop machine. In these kind of cases, you might be able to use the library out of the box without any issues. For any other scenario, please use the directions in following sections to get the library and make it work if available.</p>
<h4 id="adreno-gpu">Adreno GPU</h4>
<p>If you do not find the library on your device, check if your device has OpenCL support on the corresponding vendor's website. If your GPU does have OpenCL support, you might have to update your OS in order to get the recent updates. If you don't find the library even after OS update, you might want to email the vendor. However, if your GPU does not have OpenCL support, there isn't much we can do.</p>
<p>If you do find the library and can't use for some reason. Try to update your device drivers using the ones provided by Qualcomm. You can download them from <a href="https://developer.qualcomm.com/mobile-development/maximize-hardware/mobile-gaming-graphics-adreno/tools-and-resources">here</a>. Ensure that you are downloading the Adreno driver but not the Adreno SDK. Follow the instructions in the package and update the drivers on your device. This should resolve any issues you have.</p>
<h4 id="mali-gpu">Mali GPU</h4>
<p>Download the SDK from <a href="http://malideveloper.arm.com/develop-for-mali/sdks/mali-opencl-sdk/">here</a>. Follow the instructions in the package and build the shared library and copy it on to your device.</p>
<h4 id="kernels">Kernels</h4>
<p>I needed to do two things to achieve my target.</p>
<ul>
<li>Convert frame data provided by Android camera from NV21 format to RGBA8888 format.<br /></li>
<li>Run laplacian operator on the RGBA8888 image data.</li>
</ul>
<p>Getting a grayscale image from NV21 format is cake walk, just read the first width*height bytes and you are done. As per the format of <a href="http://www.fourcc.org/yuv.php#NV21">NV21</a>, for each 2x2 sub-matrix of pixels, there will be a pair (v,u). Since 4 pixels are reading same (v,u), i decided to use shared memory to reduce reading latency. The code below is the core of this conversion kernel, that is expected to speed up the conversion if compared to a pure C code.</p>
<pre class="sourceCode c"><code class="sourceCode c"><span class="ot">#define BLK_SIZE    16</span>
<span class="co">//prefix D for decoding</span>
<span class="ot">#define DSHRD_LEN   (BLK_SIZE/2)</span>
<span class="ot">#define DSHRD_SIZE  (2*DSHRD_LEN*DSHRD_LEN)</span>
<span class="co">// ...</span>
<span class="co">// ..</span>
<span class="co">// kernel signature</span>
{
    __local uchar uvShrd[DSHRD_SIZE];
    <span class="dt">int</span> gx  = get_global_id(<span class="dv">0</span>);
    <span class="dt">int</span> gy  = get_global_id(<span class="dv">1</span>);
    <span class="dt">int</span> lx  = get_local_id(<span class="dv">0</span>);
    <span class="dt">int</span> ly  = get_local_id(<span class="dv">1</span>);
    <span class="dt">int</span> off = im_width*im_height;
    <span class="dt">int</span> inIdx= gy*im_width+gx;
    <span class="dt">int</span> uvIdx= off + (gy/<span class="dv">2</span>)*im_width + (gx &amp; ~<span class="dv">1</span>);   <span class="co">// took me a while to figure out this formula</span>
    <span class="dt">int</span> shlx = lx/<span class="dv">2</span>;
    <span class="dt">int</span> shly = ly/<span class="dv">2</span>;
    <span class="dt">int</span> shIdx= <span class="dv">2</span>*(shlx+shly*DSHRD_LEN);
    <span class="kw">if</span>( gx%<span class="dv">2</span>==<span class="dv">0</span> &amp;&amp; gy%<span class="dv">2</span>==<span class="dv">0</span> ) {
        uvShrd[shIdx<span class="dv">+0</span>] = in[uvIdx<span class="dv">+0</span>];
        uvShrd[shIdx<span class="dv">+1</span>] = in[uvIdx<span class="dv">+1</span>];
    }
    <span class="dt">int</span> y   = (<span class="bn">0xFF</span> &amp; ((<span class="dt">int</span>)in[inIdx]));
    <span class="dt">int</span> y   = (<span class="bn">0xFF</span> &amp; ((<span class="dt">int</span>)in[inIdx]));
    <span class="kw">if</span>( y &lt; <span class="dv">16</span> ) y=<span class="dv">16</span>;
    barrier(CLK_LOCAL_MEM_FENCE);
    <span class="co">//..... rest of code that does color space conversion</span>
}</code></pre>
<p>Laplacian kernel is pretty straightforward, read all eight neighbors pixel values using shared memory and compute the convolution result. You can find kernels source code <a href="https://github.com/arrayfire/androidcl/blob/master/assets/kernels.cl">here</a>.</p>
<h4 id="auxilary-code">Auxilary Code</h4>
<p>The JNI native function declarations can be seen in the file <a href="https://github.com/arrayfire/androidcl/blob/master/jni/processor.h">processor.h</a>. I have included couple of extra Android related headers which are required to read bitmap files and print debug messages from JNI routines. The C-string pointed by macro <code>app_name</code> can be used to prune the output of the command <code>adb logcat</code> to get only messages that are associated with our JNI routines. To do that, you need to pass the value of <code>app_name</code> to adb command as an option as shown below.</p>
<pre class="sourceCode bash"><code class="sourceCode bash"><span class="kw">adb</span> logcat -s <span class="st">&quot;JNIProcessor&quot;</span></code></pre>
<p>This is quite useful when you are debugging problems that are related to accuracy of results.</p>
<p>The full source code for JNI native functions can be seen in the file <a href="https://github.com/arrayfire/androidcl/blob/master/jni/processor.cpp">processor.cpp</a>. Following are brief descriptions of the Android/JNI functions I have used in this project.</p>
<ul>
<li><code>AndroidBitmap_getInfo</code> - provides details like width, height, stride, format.<br /></li>
<li><code>AndroidBitmap_lockPixels</code> - gives us a pointer to the memory location of image data. Once lockPixels is called, it has to be followed by a unlockPixels at some point later in the code.<br /></li>
<li><code>GetPrimitiveArrayCritical</code> - as the name suggets, this is some kind of performance critical call. Hence, we should take care that the code amid a GetPrimitiveArrayCritical call and the corresponding ReleasePrimitiveArrayCritical call doesn't stall execution for extended periods of time. As per <a href="http://docs.oracle.com/javase/6/docs/technotes/guides/jni/spec/functions.html">Oracle's Java documentation</a>, a VM may temporarily disable garbage collection when the native code is holding a pointer to an array obtained via GetPrimitiveArrayCritical. So, just be careful when you use this function.<br /></li>
<li><code>ReleasePrimitiveArrayCritical</code> - used to release the memory space requested by previous GetPrimitiveArrayCritical. call.<br /></li>
<li><code>AndroidBitmap_unlockPixels</code> - used to release memory space used by a previous lockPixels call.</li>
</ul>
<p>Listed below are some of the other auxilary functions that i needed to make the code look clean.</p>
<ul>
<li><code>throwJavaException</code> - convenience function used to throw Java exceptions from JNI routines.<br /></li>
<li><code>cb</code> - callback function provided when OpenCL kernels are built. When you actually look at this, you might wonder why did i do this part alone using C-interface of OpenCL. I just wanted to try it out, hence it happened so.<br /></li>
<li><code>Java_com_example_LiveFeatureActivity_compileKernels</code> - as you might have noticed, this is a JNI routine i used to compile the OpenCL kernels only once during the app start up.</li>
</ul>
<h3 id="build-on-linux">Build (on Linux)</h3>
<p>Android-NDK provides a very useful tool <code>ndk-build</code> to build native code for Android.</p>
<ol>
<li>Clone the git repo <code>git clone https://github.com/arrayfire/androidcl</code><br /></li>
<li><code>cd androidcl</code><br /></li>
<li>Ensure android-ndk installed and environment variable <code>ANDROID_NDK</code> is pointing to ndk installation root.<br /></li>
<li>Ensure environment variable <code>PATH</code> has the path <code>$ANDROID_NDK</code> appended to it.<br /></li>
<li>Run <code>ndk-build</code> from project root.<br /></li>
<li>If you have <code>eclipse</code> installed, then import the project from eclipse and build it as usually.<br /></li>
<li>If you do not have eclipse or don't want to use it, then I assume you might already know about apache-ant tool. Run <code>ant debug</code></li>
</ol>
<h4 id="android.mk">Android.mk</h4>
<p>There are couple of things that you need to do to use OpenCL on Android apart from standard JNI makefile instructions.</p>
<ul>
<li>Ensure that you set <code>LOCAL_ARM_MODE</code> to <code>arm</code>, which enables you to build for 32-bit ARM architecture instead of the default 16-bit.<br /></li>
<li>Ensure that the variable <code>LOCAL_CPPFLAGS</code> has been set to point to OpenCL headers.<br /></li>
<li>Add <code>libOpenCL.so</code> dependency to <code>LOCAL_LDLIBS</code> variable.</li>
</ul>
<p>Click <a href="https://github.com/arrayfire/androidcl/blob/master/jni/Android.mk">here</a> to see full content of Android.mk.</p>
<p>Note: In case you want a static library, replace <code>include $(BUILD_SHARED_LIBRARY)</code> with <code>include $(BUILD_STATIC_LIBRARY)</code>.</p>
<h4 id="application.mk">Application.mk</h4>
<p>This file is not mandatory, but you may need it in cases where you want specify your requirements about the system libraries you want to use in your application. In this project i have set <code>APP_STL</code> to <code>gnustl_static</code> to declare that i want to use the static version of gnustl. I have also set the Android API target version as 14 by setting <code>APP_PLATFORM</code> to <code>android-14</code>.<br />Click <a href="https://github.com/arrayfire/androidcl/blob/master/jni/Application.mk">here</a> to see full content of Application.mk.</p>
<h3 id="app-installation">App Installation</h3>
<h4 id="ant-users">ant-users</h4>
<p>For installing the app in your mobile device do <code>ant debug install</code>. Should you choose to remove it after you are done playing with it, just do <code>ant debug uninstall</code>.</p>
<h4 id="eclipse-users">eclipse-users</h4>
<p>Right click the project in workspace and click the menu item &quot;Run As -&gt; Android Application&quot;.</p>
<h3>Video of the app running live</h3>
[video width="1280" height="720" mp4="http://sandbox.arrayfire.com/wp-content/uploads/2014/05/ocl_demo.mp4"]Screen capture of the live app[/video]