chore: import upstream snapshot with attribution

samples/android/tutorial-4-opencl/CMakeLists.txt

@@ -0,0 +1,17 @@
+set(sample example-tutorial-4-opencl)
+
+if(BUILD_FAT_JAVA_LIB)
+  set(native_deps opencv_java)
+else()
+  set(native_deps opencv_imgproc)
+endif()
+
+add_android_project(${sample} "${CMAKE_CURRENT_SOURCE_DIR}"
+    LIBRARY_DEPS "${OPENCV_ANDROID_LIB_DIR}"
+    SDK_TARGET 21 "${ANDROID_SDK_TARGET}"
+    NATIVE_DEPS ${native_deps} -lGLESv2 -lEGL
+    COPY_LIBS YES
+)
+if(TARGET ${sample})
+  add_dependencies(opencv_android_examples ${sample})
+endif()

samples/android/tutorial-4-opencl/build.gradle.in

@@ -0,0 +1,65 @@
+apply plugin: 'com.android.application'
+
+android {
+    namespace 'org.opencv.samples.tutorial4'
+    compileSdkVersion @ANDROID_COMPILE_SDK_VERSION@
+    defaultConfig {
+        applicationId "org.opencv.samples.tutorial4"
+        minSdkVersion @ANDROID_MIN_SDK_VERSION@
+        targetSdkVersion @ANDROID_TARGET_SDK_VERSION@
+        versionCode 301
+        versionName "3.01"
+
+        externalNativeBuild {
+            cmake {
+                if (gradle.opencv_source == "sdk_path") {
+                    arguments "-DOpenCV_DIR=" + project(':opencv').projectDir + "/@ANDROID_PROJECT_JNI_PATH@",
+                              "-DOPENCV_FROM_SDK=TRUE",
+                              "-DANDROID_SUPPORT_FLEXIBLE_PAGE_SIZES=ON",
+                              "-DANDROID_OPENCL_SDK=@ANDROID_OPENCL_SDK@" @OPENCV_ANDROID_CMAKE_EXTRA_ARGS@
+
+                } else {
+                    arguments "-DOPENCV_VERSION_MAJOR=@OPENCV_VERSION_MAJOR@",
+                              "-DOPENCV_FROM_SDK=FALSE",
+                              "-DANDROID_SUPPORT_FLEXIBLE_PAGE_SIZES=ON",
+                              "-DANDROID_OPENCL_SDK=@ANDROID_OPENCL_SDK@" @OPENCV_ANDROID_CMAKE_EXTRA_ARGS@
+                }
+                targets "JNIpart"
+            }
+        }
+    }
+    buildTypes {
+        release {
+            minifyEnabled false
+            proguardFiles getDefaultProguardFile('proguard-android.txt'), 'proguard-rules.pro'
+        }
+    }
+    sourceSets {
+        main {
+            java.srcDirs = @ANDROID_SAMPLE_JAVA_PATH@
+            res.srcDirs = @ANDROID_SAMPLE_RES_PATH@
+            manifest.srcFile '@ANDROID_SAMPLE_MANIFEST_PATH@'
+        }
+    }
+    externalNativeBuild {
+        cmake {
+             path '@ANDROID_SAMPLE_JNI_PATH@/CMakeLists.txt'
+        }
+    }
+    buildFeatures {
+        if (gradle.opencv_source == "maven_local" || gradle.opencv_source == "maven_central") {
+            prefab true
+        }
+    }
+}
+
+dependencies {
+    //implementation fileTree(dir: 'libs', include: ['*.jar'])
+    if (gradle.opencv_source == "sdk_path") {
+        println 'Using OpenCV from SDK'
+        implementation project(':opencv')
+    } else if (gradle.opencv_source == "maven_local" || gradle.opencv_source == "maven_central") {
+        println 'Using OpenCV from Maven repo'
+        implementation 'org.opencv:opencv:@OPENCV_VERSION_PLAIN@'
+    }
+}

samples/android/tutorial-4-opencl/gradle/AndroidManifest.xml

@@ -0,0 +1,28 @@
+<?xml version="1.0" encoding="utf-8"?>
+<manifest xmlns:android="http://schemas.android.com/apk/res/android"
+    package="org.opencv.samples.tutorial4"
+    android:versionCode="1"
+    android:versionName="1.0" >
+
+    <uses-feature android:glEsVersion="0x00020000" android:required="true"/>
+    <uses-feature android:name="android.hardware.camera"/>
+    <uses-feature android:name="android.hardware.camera2" android:required="false"/>
+    <uses-permission android:name="android.permission.CAMERA"/>
+
+    <application
+        android:allowBackup="true"
+        android:icon="@drawable/icon">
+        <activity
+            android:exported="true"
+            android:name=".Tutorial4Activity"
+            android:label="@string/app_name"
+            android:screenOrientation="landscape"
+            android:configChanges="keyboardHidden|orientation">
+            <intent-filter>
+                <action android:name="android.intent.action.MAIN" />
+                <category android:name="android.intent.category.LAUNCHER" />
+            </intent-filter>
+        </activity>
+    </application>
+
+</manifest>

samples/android/tutorial-4-opencl/jni/CLprocessor.cpp

@@ -0,0 +1,309 @@
+#ifdef OPENCL_FOUND
+#define __CL_ENABLE_EXCEPTIONS
+#define CL_USE_DEPRECATED_OPENCL_1_1_APIS /*let's give a chance for OpenCL 1.1 devices*/
+#include <CL/opencl.hpp>
+#endif
+
+#include <GLES2/gl2.h>
+#include <EGL/egl.h>
+
+#include <opencv2/core.hpp>
+#include <opencv2/imgproc.hpp>
+#include <opencv2/core/ocl.hpp>
+
+#include "common.hpp"
+#include "CLprocessor.hpp"
+
+#ifdef OPENCL_FOUND
+const char oclProgB2B[] = "// clBuffer to clBuffer";
+const char oclProgI2B[] = "// clImage to clBuffer";
+const char oclProgI2I[] = \
+  "__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; \n" \
+    "\n" \
+    "__kernel void Laplacian( \n" \
+    "        __read_only image2d_t imgIn, \n" \
+    "        __write_only image2d_t imgOut \n" \
+    "    ) { \n" \
+    "  \n" \
+    "    const int2 pos = {get_global_id(0), get_global_id(1)}; \n" \
+    "  \n" \
+    "    float4 sum = (float4) 0.0f; \n" \
+    "    sum += read_imagef(imgIn, sampler, pos + (int2)(-1,0)); \n" \
+    "    sum += read_imagef(imgIn, sampler, pos + (int2)(+1,0)); \n" \
+    "    sum += read_imagef(imgIn, sampler, pos + (int2)(0,-1)); \n" \
+    "    sum += read_imagef(imgIn, sampler, pos + (int2)(0,+1)); \n" \
+    "    sum -= read_imagef(imgIn, sampler, pos) * 4; \n" \
+    "  \n" \
+    "    write_imagef(imgOut, pos, sum*10); \n" \
+    "} \n";
+
+static void dumpCLinfo()
+{
+    LOGD("*** OpenCL info ***");
+    try
+    {
+        std::vector<cl::Platform> platforms;
+        cl::Platform::get(&platforms);
+        LOGD("OpenCL info: Found %d OpenCL platforms", platforms.size());
+        for (int i = 0; i < platforms.size(); ++i)
+        {
+            std::string name = platforms[i].getInfo<CL_PLATFORM_NAME>();
+            std::string version = platforms[i].getInfo<CL_PLATFORM_VERSION>();
+            std::string profile = platforms[i].getInfo<CL_PLATFORM_PROFILE>();
+            std::string extensions = platforms[i].getInfo<CL_PLATFORM_EXTENSIONS>();
+            LOGD( "OpenCL info: Platform[%d] = %s, ver = %s, prof = %s, ext = %s",
+                  i, name.c_str(), version.c_str(), profile.c_str(), extensions.c_str() );
+        }
+
+        std::vector<cl::Device> devices;
+        platforms[0].getDevices(CL_DEVICE_TYPE_ALL, &devices);
+
+        for (int i = 0; i < devices.size(); ++i)
+        {
+            std::string name = devices[i].getInfo<CL_DEVICE_NAME>();
+            std::string extensions = devices[i].getInfo<CL_DEVICE_EXTENSIONS>();
+            cl_ulong type = devices[i].getInfo<CL_DEVICE_TYPE>();
+            LOGD( "OpenCL info: Device[%d] = %s (%s), ext = %s",
+                  i, name.c_str(), (type==CL_DEVICE_TYPE_GPU ? "GPU" : "CPU"), extensions.c_str() );
+        }
+    }
+    catch(const cl::Error& e)
+    {
+        LOGE( "OpenCL info: error while gathering OpenCL info: %s (%d)", e.what(), e.err() );
+    }
+    catch(const std::exception& e)
+    {
+        LOGE( "OpenCL info: error while gathering OpenCL info: %s", e.what() );
+    }
+    catch(...)
+    {
+        LOGE( "OpenCL info: unknown error while gathering OpenCL info" );
+    }
+    LOGD("*******************");
+}
+
+cl::Context theContext;
+cl::CommandQueue theQueue;
+cl::Program theProgB2B, theProgI2B, theProgI2I;
+bool haveOpenCL = false;
+
+//![init_opencl]
+int initCL()
+{
+    dumpCLinfo();
+    LOGE("initCL: start initCL");
+    EGLDisplay mEglDisplay = eglGetCurrentDisplay();
+    if (mEglDisplay == EGL_NO_DISPLAY)
+        LOGE("initCL: eglGetCurrentDisplay() returned 'EGL_NO_DISPLAY', error = %x", eglGetError());
+
+    EGLContext mEglContext = eglGetCurrentContext();
+    if (mEglContext == EGL_NO_CONTEXT)
+        LOGE("initCL: eglGetCurrentContext() returned 'EGL_NO_CONTEXT', error = %x", eglGetError());
+
+    cl_context_properties props[] =
+    {   CL_GL_CONTEXT_KHR,   (cl_context_properties) mEglContext,
+        CL_EGL_DISPLAY_KHR,  (cl_context_properties) mEglDisplay,
+        CL_CONTEXT_PLATFORM, 0,
+        0 };
+
+    try
+    {
+        haveOpenCL = false;
+        cl::Platform p = cl::Platform::getDefault();
+        std::string ext = p.getInfo<CL_PLATFORM_EXTENSIONS>();
+        if(ext.find("cl_khr_gl_sharing") == std::string::npos)
+            LOGE("Warning: CL-GL sharing isn't supported by PLATFORM");
+        props[5] = (cl_context_properties) p();
+
+        theContext = cl::Context(CL_DEVICE_TYPE_GPU, props);
+        std::vector<cl::Device> devs = theContext.getInfo<CL_CONTEXT_DEVICES>();
+        LOGD("Context returned %d devices, taking the 1st one", devs.size());
+        ext = devs[0].getInfo<CL_DEVICE_EXTENSIONS>();
+        if(ext.find("cl_khr_gl_sharing") == std::string::npos)
+            LOGE("Warning: CL-GL sharing isn't supported by DEVICE");
+
+        theQueue = cl::CommandQueue(theContext, devs[0]);
+
+        cl::Program::Sources src(1, std::make_pair(oclProgI2I, sizeof(oclProgI2I)));
+        theProgI2I = cl::Program(theContext, src);
+        theProgI2I.build(devs);
+
+        cv::ocl::attachContext(p.getInfo<CL_PLATFORM_NAME>(), p(), theContext(), devs[0]());
+        if( cv::ocl::useOpenCL() )
+            LOGD("OpenCV+OpenCL works OK!");
+        else
+            LOGE("Can't init OpenCV with OpenCL TAPI");
+        haveOpenCL = true;
+    }
+    catch(const cl::Error& e)
+    {
+        LOGE("cl::Error: %s (%d)", e.what(), e.err());
+        return 1;
+    }
+    catch(const std::exception& e)
+    {
+        LOGE("std::exception: %s", e.what());
+        return 2;
+    }
+    catch(...)
+    {
+        LOGE( "OpenCL info: unknown error while initializing OpenCL stuff" );
+        return 3;
+    }
+    LOGD("initCL completed");
+
+    if (haveOpenCL)
+        return 0;
+    else
+        return 4;
+}
+//![init_opencl]
+
+#define GL_TEXTURE_2D 0x0DE1
+void procOCL_I2I(int texIn, int texOut, int w, int h)
+{
+    LOGD("Processing OpenCL Direct (image2d)");
+    if(!haveOpenCL)
+    {
+        LOGE("OpenCL isn't initialized");
+        return;
+    }
+
+    LOGD("procOCL_I2I(%d, %d, %d, %d)", texIn, texOut, w, h);
+//![process_pure_opencl]
+    cl::ImageGL imgIn (theContext, CL_MEM_READ_ONLY,  GL_TEXTURE_2D, 0, texIn);
+    cl::ImageGL imgOut(theContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, texOut);
+    std::vector < cl::Memory > images;
+    images.push_back(imgIn);
+    images.push_back(imgOut);
+
+    int64_t t = getTimeMs();
+    theQueue.enqueueAcquireGLObjects(&images);
+    theQueue.finish();
+    LOGD("enqueueAcquireGLObjects() costs %d ms", getTimeInterval(t));
+
+    t = getTimeMs();
+    cl::Kernel Laplacian(theProgI2I, "Laplacian"); //TODO: may be done once
+    Laplacian.setArg(0, imgIn);
+    Laplacian.setArg(1, imgOut);
+    theQueue.finish();
+    LOGD("Kernel() costs %d ms", getTimeInterval(t));
+
+    t = getTimeMs();
+    theQueue.enqueueNDRangeKernel(Laplacian, cl::NullRange, cl::NDRange(w, h), cl::NullRange);
+    theQueue.finish();
+    LOGD("enqueueNDRangeKernel() costs %d ms", getTimeInterval(t));
+
+    t = getTimeMs();
+    theQueue.enqueueReleaseGLObjects(&images);
+    theQueue.finish();
+    LOGD("enqueueReleaseGLObjects() costs %d ms", getTimeInterval(t));
+//![process_pure_opencl]
+}
+
+void procOCL_OCV(int texIn, int texOut, int w, int h)
+{
+    LOGD("Processing OpenCL via OpenCV");
+    if(!haveOpenCL)
+    {
+        LOGE("OpenCL isn't initialized");
+        return;
+    }
+
+//![process_tapi]
+    int64_t t = getTimeMs();
+    cl::ImageGL imgIn (theContext, CL_MEM_READ_ONLY,  GL_TEXTURE_2D, 0, texIn);
+    std::vector < cl::Memory > images(1, imgIn);
+    theQueue.enqueueAcquireGLObjects(&images);
+    theQueue.finish();
+    cv::UMat uIn, uOut, uTmp;
+    cv::ocl::convertFromImage(imgIn(), uIn);
+    LOGD("loading texture data to OpenCV UMat costs %d ms", getTimeInterval(t));
+    theQueue.enqueueReleaseGLObjects(&images);
+
+    t = getTimeMs();
+    //cv::blur(uIn, uOut, cv::Size(5, 5));
+    cv::Laplacian(uIn, uTmp, CV_8U);
+    cv:multiply(uTmp, 10, uOut);
+    cv::ocl::finish();
+    LOGD("OpenCV processing costs %d ms", getTimeInterval(t));
+
+    t = getTimeMs();
+    cl::ImageGL imgOut(theContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, texOut);
+    images.clear();
+    images.push_back(imgOut);
+    theQueue.enqueueAcquireGLObjects(&images);
+    cl_mem clBuffer = (cl_mem)uOut.handle(cv::ACCESS_READ);
+    cl_command_queue q = (cl_command_queue)cv::ocl::Queue::getDefault().ptr();
+    size_t offset = 0;
+    size_t origin[3] = { 0, 0, 0 };
+    size_t region[3] = { (size_t)w, (size_t)h, 1 };
+    CV_Assert(clEnqueueCopyBufferToImage (q, clBuffer, imgOut(), offset, origin, region, 0, NULL, NULL) == CL_SUCCESS);
+    theQueue.enqueueReleaseGLObjects(&images);
+    cv::ocl::finish();
+    LOGD("uploading results to texture costs %d ms", getTimeInterval(t));
+//![process_tapi]
+}
+#else
+int initCL()
+{
+    return 5;
+}
+#endif
+
+void closeCL()
+{
+}
+
+void drawFrameProcCPU(int w, int h, int texOut)
+{
+    LOGD("Processing on CPU");
+    int64_t t;
+
+    // let's modify pixels in FBO texture in C++ code (on CPU)
+    static cv::Mat m;
+    m.create(h, w, CV_8UC4);
+
+    // read
+    t = getTimeMs();
+    // expecting FBO to be bound
+    glReadPixels(0, 0, w, h, GL_RGBA, GL_UNSIGNED_BYTE, m.data);
+    LOGD("glReadPixels() costs %d ms", getTimeInterval(t));
+
+   // modify
+    t = getTimeMs();
+    cv::Laplacian(m, m, CV_8U);
+    m *= 10;
+    LOGD("Laplacian() costs %d ms", getTimeInterval(t));
+
+    // write back
+    glActiveTexture(GL_TEXTURE0);
+    glBindTexture(GL_TEXTURE_2D, texOut);
+    t = getTimeMs();
+    glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, w, h, GL_RGBA, GL_UNSIGNED_BYTE, m.data);
+    LOGD("glTexSubImage2D() costs %d ms", getTimeInterval(t));
+}
+
+
+enum ProcMode {PROC_MODE_NO_PROC=0, PROC_MODE_CPU=1, PROC_MODE_OCL_DIRECT=2, PROC_MODE_OCL_OCV=3};
+
+void processFrame(int tex1, int tex2, int w, int h, int mode)
+{
+    switch(mode)
+    {
+        //case PROC_MODE_NO_PROC:
+    case PROC_MODE_CPU:
+        drawFrameProcCPU(w, h, tex2);
+        break;
+#ifdef OPENCL_FOUND
+    case PROC_MODE_OCL_DIRECT:
+        procOCL_I2I(tex1, tex2, w, h);
+        break;
+    case PROC_MODE_OCL_OCV:
+        procOCL_OCV(tex1, tex2, w, h);
+        break;
+#endif
+    default:
+        LOGE("Unexpected processing mode: %d", mode);
+    }
+}

samples/android/tutorial-4-opencl/jni/CLprocessor.hpp

@@ -0,0 +1,8 @@
+#ifndef __CL_PROCESSOR_HPP__
+#define __CL_PROCESSOR_HPP__
+
+int initCL();
+void closeCL();
+void processFrame(int tex1, int tex2, int w, int h, int mode);
+
+#endif

samples/android/tutorial-4-opencl/jni/CMakeLists.txt

@@ -0,0 +1,50 @@
+cmake_minimum_required(VERSION 3.6)
+
+set(target JNIpart)
+project(${target} CXX)
+
+if (OPENCV_FROM_SDK)
+  message(STATUS "Using OpenCV from local SDK")
+  set(ANDROID_OPENCV_COMPONENTS "opencv_java" CACHE STRING "")
+else()
+  message(STATUS "Using OpenCV from AAR (Maven repo)")
+  set(ANDROID_OPENCV_COMPONENTS "OpenCV::opencv_java${OPENCV_VERSION_MAJOR}" CACHE STRING "")
+endif()
+
+message(STATUS "ANDROID_ABI=${ANDROID_ABI}")
+find_package(OpenCV REQUIRED COMPONENTS ${ANDROID_OPENCV_COMPONENTS})
+find_package(OpenCL QUIET)
+
+file(GLOB srcs *.cpp *.c)
+file(GLOB hdrs *.hpp *.h)
+
+# For 16k pages support with NDK prior 27
+# Details: https://developer.android.com/guide/practices/page-sizes?hl=en
+if(ANDROID_SUPPORT_FLEXIBLE_PAGE_SIZES)
+  if(ANDROID_ABI STREQUAL arm64-v8a OR ANDROID_ABI STREQUAL x86_64)
+    set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -Wl,-z,max-page-size=16384")
+  endif()
+endif()
+
+include_directories("${CMAKE_CURRENT_LIST_DIR}")
+add_library(${target} SHARED ${srcs} ${hdrs})
+
+target_link_libraries(${target} ${ANDROID_OPENCV_COMPONENTS} -lGLESv2 -lEGL -llog)
+
+if(OpenCL_FOUND)
+  include_directories(${OpenCL_INCLUDE_DIRS})
+  target_link_libraries(${target} ${OpenCL_LIBRARIES})
+  add_definitions("-DOPENCL_FOUND")
+elseif(NOT ("${ANDROID_OPENCL_SDK}" STREQUAL ""))
+  include_directories(${ANDROID_OPENCL_SDK}/include)
+  link_directories(${ANDROID_OPENCL_SDK}/lib)
+  target_link_directories(${target} PRIVATE ${ANDROID_OPENCL_SDK}/lib)
+
+  set_target_properties(${target} PROPERTIES LINK_FLAGS "-Wl,--allow-shlib-undefined")
+  target_link_libraries(${target} -lOpenCL)
+
+  add_definitions("-DOPENCL_FOUND")
+  add_definitions("-DCL_HPP_MINIMUM_OPENCL_VERSION=120")
+  add_definitions("-DCL_HPP_TARGET_OPENCL_VERSION=120")
+  add_definitions("-DCL_HPP_ENABLE_PROGRAM_CONSTRUCTION_FROM_ARRAY_COMPATIBILITY")
+endif()

samples/android/tutorial-4-opencl/jni/common.hpp

@@ -0,0 +1,19 @@
+#include <android/log.h>
+#define LOG_TAG "JNIpart"
+//#define LOGD(...)
+#define LOGD(...) ((void)__android_log_print(ANDROID_LOG_DEBUG, LOG_TAG, __VA_ARGS__))
+#define LOGE(...) ((void)__android_log_print(ANDROID_LOG_ERROR, LOG_TAG, __VA_ARGS__))
+
+#include <time.h> // clock_gettime
+
+static inline int64_t getTimeMs()
+{
+    struct timespec now;
+    clock_gettime(CLOCK_MONOTONIC, &now);
+    return (int64_t) now.tv_sec*1000 + now.tv_nsec/1000000;
+}
+
+static inline int getTimeInterval(int64_t startTime)
+{
+    return int(getTimeMs() - startTime);
+}

samples/android/tutorial-4-opencl/jni/jni.cpp

@@ -0,0 +1,36 @@
+#include <jni.h>
+#include "CLprocessor.hpp"
+
+extern "C" {
+JNIEXPORT jboolean JNICALL Java_org_opencv_samples_tutorial4_NativePart_builtWithOpenCL(JNIEnv * env, jclass cls);
+
+JNIEXPORT jint JNICALL Java_org_opencv_samples_tutorial4_NativePart_initCL(JNIEnv * env, jclass cls);
+
+JNIEXPORT void JNICALL Java_org_opencv_samples_tutorial4_NativePart_closeCL(JNIEnv * env, jclass cls);
+
+JNIEXPORT void JNICALL Java_org_opencv_samples_tutorial4_NativePart_processFrame(JNIEnv * env, jclass cls, jint tex1, jint tex2, jint w, jint h, jint mode);
+
+JNIEXPORT jboolean JNICALL Java_org_opencv_samples_tutorial4_NativePart_builtWithOpenCL(JNIEnv * env, jclass cls)
+{
+#ifdef OPENCL_FOUND
+    return JNI_TRUE;
+#else
+    return JNI_FALSE;
+#endif
+}
+
+JNIEXPORT jint JNICALL Java_org_opencv_samples_tutorial4_NativePart_initCL(JNIEnv * env, jclass cls)
+{
+    return initCL();
+}
+
+JNIEXPORT void JNICALL Java_org_opencv_samples_tutorial4_NativePart_closeCL(JNIEnv * env, jclass cls)
+{
+    closeCL();
+}
+
+JNIEXPORT void JNICALL Java_org_opencv_samples_tutorial4_NativePart_processFrame(JNIEnv * env, jclass cls, jint tex1, jint tex2, jint w, jint h, jint mode)
+{
+    processFrame(tex1, tex2, w, h, mode);
+}
+} // extern "C"

samples/android/tutorial-4-opencl/lint.xml

@@ -0,0 +1,3 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<lint>
+</lint>

samples/android/tutorial-4-opencl/res/layout/activity.xml

@@ -0,0 +1,26 @@
+<FrameLayout xmlns:android="http://schemas.android.com/apk/res/android"
+    android:layout_width="match_parent"
+    android:layout_height="match_parent" >
+
+    <org.opencv.samples.tutorial4.MyGLSurfaceView
+        android:layout_width="match_parent"
+        android:layout_height="match_parent"
+        android:id="@+id/my_gl_surface_view" />
+
+    <LinearLayout
+            android:layout_width="match_parent"
+            android:layout_height="match_parent"
+            android:orientation = "vertical" >
+	    <TextView
+	        android:layout_width="wrap_content"
+	        android:layout_height="wrap_content"
+	        android:id="@+id/fps_text_view"
+	        android:text="FPS:" />
+        <TextView
+            android:layout_width="wrap_content"
+            android:layout_height="wrap_content"
+            android:id="@+id/proc_mode_text_view"
+            android:text="Processing mode:" />
+	    </LinearLayout>
+
+</FrameLayout>

samples/android/tutorial-4-opencl/res/values/strings.xml

@@ -0,0 +1,6 @@
+<?xml version="1.0" encoding="utf-8"?>
+<resources>
+
+    <string name="app_name">OpenCV Tutorial 4 - Use OpenCL</string>
+
+</resources>

samples/android/tutorial-4-opencl/src/org/opencv/samples/tutorial4/MyGLSurfaceView.java

@@ -0,0 +1,115 @@
+package org.opencv.samples.tutorial4;
+
+import org.opencv.android.CameraGLSurfaceView;
+
+import android.app.Activity;
+import android.content.Context;
+import android.os.Handler;
+import android.os.Looper;
+import android.util.AttributeSet;
+import android.util.Log;
+import android.view.MotionEvent;
+import android.view.SurfaceHolder;
+import android.widget.TextView;
+import android.widget.Toast;
+
+//![minimal_surface_view]
+public class MyGLSurfaceView extends CameraGLSurfaceView implements CameraGLSurfaceView.CameraTextureListener {
+
+    static final String LOGTAG = "MyGLSurfaceView";
+    protected int procMode = NativePart.PROCESSING_MODE_NO_PROCESSING;
+    static final String[] procModeName = new String[] {"No Processing", "CPU", "OpenCL Direct", "OpenCL via OpenCV"};
+    protected int  frameCounter;
+    protected long lastNanoTime;
+    TextView mFpsText = null;
+
+    public MyGLSurfaceView(Context context, AttributeSet attrs) {
+        super(context, attrs);
+    }
+
+    @Override
+    public boolean onTouchEvent(MotionEvent e) {
+        if(e.getAction() == MotionEvent.ACTION_DOWN)
+            ((Activity)getContext()).openOptionsMenu();
+        return true;
+    }
+
+    @Override
+    public void surfaceCreated(SurfaceHolder holder) {
+        super.surfaceCreated(holder);
+        //NativePart.initCL();
+    }
+
+    @Override
+    public void surfaceDestroyed(SurfaceHolder holder) {
+        //NativePart.closeCL();
+        super.surfaceDestroyed(holder);
+    }
+
+    public void setProcessingMode(int newMode) {
+        if(newMode>=0 && newMode<procModeName.length)
+            procMode = newMode;
+        else
+            Log.e(LOGTAG, "Ignoring invalid processing mode: " + newMode);
+
+        ((Activity) getContext()).runOnUiThread(new Runnable() {
+            public void run() {
+                Toast.makeText(getContext(), "Selected mode: " + procModeName[procMode], Toast.LENGTH_LONG).show();
+            }
+        });
+    }
+
+    @Override
+    public void onCameraViewStarted(int width, int height) {
+        ((Activity) getContext()).runOnUiThread(new Runnable() {
+            public void run() {
+                Toast.makeText(getContext(), "onCameraViewStarted", Toast.LENGTH_SHORT).show();
+            }
+        });
+        if (NativePart.builtWithOpenCL())
+            NativePart.initCL();
+        frameCounter = 0;
+        lastNanoTime = System.nanoTime();
+    }
+
+    @Override
+    public void onCameraViewStopped() {
+        ((Activity) getContext()).runOnUiThread(new Runnable() {
+            public void run() {
+                Toast.makeText(getContext(), "onCameraViewStopped", Toast.LENGTH_SHORT).show();
+            }
+        });
+    }
+
+    @Override
+    public boolean onCameraTexture(int texIn, int texOut, int width, int height) {
+        // FPS
+        frameCounter++;
+        if(frameCounter >= 30)
+        {
+            final int fps = (int) (frameCounter * 1e9 / (System.nanoTime() - lastNanoTime));
+            Log.i(LOGTAG, "drawFrame() FPS: "+fps);
+            if(mFpsText != null) {
+                Runnable fpsUpdater = new Runnable() {
+                    public void run() {
+                        mFpsText.setText("FPS: " + fps);
+                    }
+                };
+                new Handler(Looper.getMainLooper()).post(fpsUpdater);
+            } else {
+                Log.d(LOGTAG, "mFpsText == null");
+                mFpsText = (TextView)((Activity) getContext()).findViewById(R.id.fps_text_view);
+            }
+            frameCounter = 0;
+            lastNanoTime = System.nanoTime();
+        }
+
+
+        if(procMode == NativePart.PROCESSING_MODE_NO_PROCESSING)
+            return false;
+
+        NativePart.processFrame(texIn, texOut, width, height, procMode);
+        return true;
+    }
+}
+//![minimal_surface_view]

samples/android/tutorial-4-opencl/src/org/opencv/samples/tutorial4/NativePart.java

@@ -0,0 +1,21 @@
+package org.opencv.samples.tutorial4;
+
+//![native_part]
+public class NativePart {
+    static
+    {
+        System.loadLibrary("opencv_java4");
+        System.loadLibrary("JNIpart");
+    }
+
+    public static final int PROCESSING_MODE_NO_PROCESSING = 0;
+    public static final int PROCESSING_MODE_CPU = 1;
+    public static final int PROCESSING_MODE_OCL_DIRECT = 2;
+    public static final int PROCESSING_MODE_OCL_OCV = 3;
+
+    public static native boolean builtWithOpenCL();
+    public static native int initCL();
+    public static native void closeCL();
+    public static native void processFrame(int tex1, int tex2, int w, int h, int mode);
+}
+//![native_part]

samples/android/tutorial-4-opencl/src/org/opencv/samples/tutorial4/Tutorial4Activity.java

@@ -0,0 +1,100 @@
+package org.opencv.samples.tutorial4;
+
+import android.content.pm.ActivityInfo;
+import android.os.Bundle;
+import android.view.Menu;
+import android.view.MenuInflater;
+import android.view.MenuItem;
+import android.view.Window;
+import android.view.WindowManager;
+import android.widget.TextView;
+
+import org.opencv.android.CameraActivity;
+
+public class Tutorial4Activity extends CameraActivity {
+
+    private MyGLSurfaceView mView;
+    private TextView mProcMode;
+
+    private boolean builtWithOpenCL = false;
+
+    private MenuItem mItemNoProc;
+    private MenuItem mItemCpu;
+    private MenuItem mItemOclDirect;
+    private MenuItem mItemOclOpenCV;
+
+    @Override
+    public void onCreate(Bundle savedInstanceState) {
+        super.onCreate(savedInstanceState);
+        requestWindowFeature(Window.FEATURE_NO_TITLE);
+        getWindow().setFlags(WindowManager.LayoutParams.FLAG_FULLSCREEN,
+                WindowManager.LayoutParams.FLAG_FULLSCREEN);
+        getWindow().setFlags(WindowManager.LayoutParams.FLAG_KEEP_SCREEN_ON,
+                WindowManager.LayoutParams.FLAG_KEEP_SCREEN_ON);
+        setRequestedOrientation(ActivityInfo.SCREEN_ORIENTATION_LANDSCAPE);
+
+        //mView = new MyGLSurfaceView(this, null);
+        //setContentView(mView);
+        setContentView(R.layout.activity);
+        mView = (MyGLSurfaceView) findViewById(R.id.my_gl_surface_view);
+        mView.setCameraTextureListener(mView);
+        TextView tv = (TextView)findViewById(R.id.fps_text_view);
+        mProcMode = (TextView)findViewById(R.id.proc_mode_text_view);
+        runOnUiThread(new Runnable() {
+            public void run() {
+                mProcMode.setText("Processing mode: No processing");
+            }
+        });
+
+        builtWithOpenCL = NativePart.builtWithOpenCL();
+        mView.setProcessingMode(NativePart.PROCESSING_MODE_NO_PROCESSING);
+    }
+
+    @Override
+    protected void onPause() {
+        mView.onPause();
+        super.onPause();
+    }
+
+    @Override
+    protected void onResume() {
+        super.onResume();
+        mView.onResume();
+    }
+
+    @Override
+    public boolean onCreateOptionsMenu(Menu menu) {
+        mItemNoProc = menu.add("No processing");
+        mItemCpu = menu.add("Use CPU code");
+        if (builtWithOpenCL) {
+            mItemOclOpenCV = menu.add("Use OpenCL via OpenCV");
+            mItemOclDirect = menu.add("Use OpenCL direct");
+        }
+        return true;
+    }
+
+    @Override
+    public boolean onOptionsItemSelected(MenuItem item) {
+        String procName = "Not selected";
+        int procMode = NativePart.PROCESSING_MODE_NO_PROCESSING;
+
+        if (item == mItemNoProc) {
+            procMode = NativePart.PROCESSING_MODE_NO_PROCESSING;
+            procName = "Processing mode: No Processing";
+        } else if (item == mItemCpu) {
+            procMode = NativePart.PROCESSING_MODE_CPU;
+            procName = "Processing mode: CPU";
+        } else if (item == mItemOclOpenCV && builtWithOpenCL) {
+            procMode = NativePart.PROCESSING_MODE_OCL_OCV;
+            procName = "Processing mode: OpenCL via OpenCV (TAPI)";
+        } else if (item == mItemOclDirect && builtWithOpenCL) {
+            procMode = NativePart.PROCESSING_MODE_OCL_DIRECT;
+            procName = "Processing mode: OpenCL direct";
+        }
+
+        mView.setProcessingMode(procMode);
+        mProcMode.setText(procName);
+
+        return true;
+    }
+}