-수정 필- [Android Studio, JNI] java 환경에서 C언어 활용(bitmap, cross compile)

0. bitmap 분석

이번 문제의 목표는 java 환경의 bitmap data를 C언어에 적용하여 필터 연산을 수행하는 것이다.

이때, 이번 실습에서 활용하는 필터는 gaussian blur 필터이다.

적용 방법은 2가지로

1. CPU 기반

2. GPU 기반

이다.

이때, 우리는 OpenCL을 활용하여 device에서 다른 장치를 활용하는 것이다.

최종적으로 CPU, GPU의 연산 속도의 차이를 비교할 것이다.

- 구조

1. 풀이

1) OpenCL을 활용할 수 있게 C언어 환경을 사용할 수 있는 Native C++ 선택

2) openCL에서 필요한 Header file 추가

C:\Users\"사용자명"\AppData\Local\Android\Sdk\ndk\21.3.6528147\toolchains\llvm\prebuilt\windows-x86_64\sysroot\usr\include\c++

+) 처음은 include에 CL header 파일을 복사하였지만 c언어 코딩을 하며 header 파일이 로드되지 않는 것을 볼 수 있었다.

 현재 코딩하는 환경인 C++환경에 파일을 복사하여 #include가 동작하는 모습을 볼 수 있었다.

- 실제 복사한 결과

- include 된 위치

- include가 문제없이 수행된 것을 볼 수 있다.

3) libGLES_mali.so 추가

CMakeList에 OpenDriver가 활용하는 라이브러리에 libGLES_mali.so를 추가한다.

이 추가를 통해 우리는 openCL 라이브러리를 활용할 수 있게 되는 것이다.

ninja: 

error: 'C:/Users/"사용자"/AndroidStudioProjects/week12_openCL_with_Cpp/app/build/OpenCL_lib_and_inlclude/libGLES_mali.so', needed by 

'C:/Users/"사용자"/AndroidStudioProjects/week12_openCL_with_Cpp/app/build/intermediates/cmake/debug/obj/armeabi-v7a/libOpenCLDriver.so', missing and no known rule to make it

1) 

2)

3)

2. 소스코드

- CMakeList.txt

# For more information about using CMake with Android Studio, read the
# documentation: https://d.android.com/studio/projects/add-native-code.html
 
# Sets the minimum version of CMake required to build the native library.
 
cmake_minimum_required(VERSION 3.10.2)
 
# Declares and names the project.
 
project("week12_opencl_with_cpp")
 
# Creates and names a library, sets it as either STATIC
# or SHARED, and provides the relative paths to its source code.
# You can define multiple libraries, and CMake builds them for you.
# Gradle automatically packages shared libraries with your APK.
 
add_library( # Sets the name of the library.
             native-lib
 
             # Sets the library as a shared library.
             SHARED
 
             # Provides a relative path to your source file(s).
             native-lib.cpp )
add_library(GLES_mali STATIC IMPORTED)
 
add_library( # Sets the name of the library.
        OpenCLDriver
 
        # Sets the library as a shared library.
        SHARED
 
        # Provides a relative path to your source file(s).
        OpenCLDriver.c )
 
# Searches for a specified prebuilt library and stores the path as a
# variable. Because CMake includes system libraries in the search path by
# default, you only need to specify the name of the public NDK library
# you want to add. CMake verifies that the library exists before
# completing its build.
 
find_library( # Sets the name of the path variable.
        log-lib
 
        # Specifies the name of the NDK library that
        # you want CMake to locate.
        log )
 
find_library( # Sets the name of the path variable.
        jnigraphics-lib
 
        # Specifies the name of the NDK library that
        # you want CMake to locate.
        jnigraphics )
 
# Specifies libraries CMake should link to your target library. You
# can link multiple libraries, such as libraries you define in this
# build script, prebuilt third-party libraries, or system libraries.
 
target_link_libraries( # Specifies the target library.
        native-lib
 
        # Links the target library to the log library
        # included in the NDK.
        ${log-lib} )
 
target_link_libraries( # Specifies the target library.
        OpenCLDriver
 
        # Links the target library to the log library
        # included in the NDK.
        C:/Users/"사용자명"/AndroidStudioProjects/week12_openCL_with_Cpp/app/build/OpenCL_lib_and_inlclude/lib/libGLES_mali.so ${log-lib} ${jnigraphics-lib} )
 
target_include_directories(native-lib PRIVATE src/main.cpp)

- OpenCLDriver.c

ver. 1

#include<jni.h>
#include<stdio.h>
#include<stdlib.h>
#include <sys/time.h>
#include <android/log.h>
#include <android/bitmap.h>
 
#include <CL/opencl.h>
#include <assert.h>
 
#define LOG_TAG "DEBUG"
#define LOGD(...) __android_log_print(ANDROID_LOG_DEBUG,LOG_TAG,__VA_ARGS__)
#define LOGE(...) __android_log_print(ANDROID_LOG_ERROR,LOG_TAG,__VA_ARGS__)
 
#define RGB8888_A(p) ((p & (0xff<<24))      >> 24 )
#define RGB8888_B(p) ((p & (0xff << 16)) >> 16 )
#define RGB8888_G(p) ((p & (0xff << 8))  >> 8 )
#define RGB8888_R(p) (p & (0xff) )
 
JNIEXPORT jobject JNICALL
Java_com_example_week12_1opencl_1with_1cpp_MainActivity_GaussianBlurBitmap
        (JNIEnv *env, jclass class , jobject bitmap)
{
    AndroidBitmapInfo bitmapInfo;
//    memset(&bitmapInfo , 0 , sizeof(bitmapInfo));
    int ret = AndroidBitmap_getInfo(env, bitmap, &bitmapInfo);
    if (ANDROID_BITMAP_RESULT_SUCCESS != ret) {
        LOGE("AndroidBitmap_getInfo() bitmap failed ! error=%d", ret)
    }
    // Get Bitmap pixel cache pointer: traverse to read BGRA data from Bitmap memory addrPtr
    void *bitmapPixels;
    ret = AndroidBitmap_lockPixels(env, bitmap, &bitmapPixels);
    if (ANDROID_BITMAP_RESULT_SUCCESS != ret) {
        LOGE("AndroidBitmap_lockPixels() bitmap failed ! error=%d", ret)
    }
 
    // Logic for performing image manipulation
    // Get width and height
    uint32_t mWidth = bitmapInfo.width;
    uint32_t mHeight = bitmapInfo.height;
    // Get native data
    //auto pixelArr = ((uint32_t *) bitmapPixels);
 
    LOGE("bitmap width = %d", mWidth)
    LOGE("bitmap height = %d", mHeight)
    LOGE("bitmap format = %d", bitmapInfo.format)
    int a,r, g, b;
 
    //
    float red = 0, green = 0, blue = 0;
    int row = 0, col = 0;
    int m, n, k,x,y;
    int pix;
    float mask[9][9] = {
            {0.011237, 0.011637, 0.011931, 0.012111, 0.012172, 0.012111, 0.011931, 0.011637, 0.011237},
            {0.011637, 0.012051, 0.012356, 0.012542, 0.012605, 0.012542, 0.012356, 0.012051, 0.011637},
            {0.011931, 0.012356, 0.012668, 0.012860, 0.012924, 0.012860, 0.012668, 0.012356, 0.011931},
            {0.012111, 0.012542, 0.012860, 0.013054, 0.013119, 0.013054, 0.012860, 0.012542, 0.012111},
            {0.012172, 0.012605, 0.012924, 0.013119, 0.013185, 0.013119, 0.012924, 0.012605, 0.012172},
            {0.012111, 0.012542, 0.012860, 0.013054, 0.013119, 0.013054, 0.012860, 0.012542, 0.012111},
            {0.011931, 0.012356, 0.012668, 0.012860, 0.012924, 0.012860, 0.012668, 0.012356, 0.011931},
            {0.011637, 0.012051, 0.012356, 0.012542, 0.012605, 0.012542, 0.012356, 0.012051, 0.011637},
            {0.011237, 0.011637, 0.011931, 0.012111, 0.012172, 0.012111, 0.011931, 0.011637, 0.011237},
    };
    //
    ////
    uint32_t* pixelArr = (uint32_t*) bitmapPixels;
    uint32_t* tempPixels = (uint32_t*)malloc(sizeof(uint32_t)*info.height * info.width);
    int stride = info.stride;
    int pixelsCount = info.height * info.width;
    memcpy(tempPixels, pixelArr, sizeof(uint32_t) * pixelsCount);
    AndroidBitmap_unlockPixels(env, bitmap);
    //
    //recycle bitmap - using bitmap.recycle()
    //
    LOGD("recycling bitmap...");
    jclass bitmapCls = env->GetObjectClass(bitmap);
    jmethodID recycleFunction = env->GetMethodID(bitmapCls, "recycle", "()V");
    if (recycleFunction == 0)
    {
        LOGE("error recycling!");
        return NULL;
    }
    env->CallVoidMethod(bitmap, recycleFunction);
    //
    //creating a new bitmap to put the pixels into it - using Bitmap Bitmap.createBitmap (int width, int height, Bitmap.Config config) :
    //
    LOGD("creating new bitmap...");
    jmethodID createBitmapFunction = env->GetStaticMethodID(bitmapCls, "createBitmap", "(IILandroid/graphics/Bitmap$Config;)Landroid/graphics/Bitmap;");
    jstring configName = env->NewStringUTF("ARGB_8888");
    jclass bitmapConfigClass = env->FindClass("android/graphics/Bitmap$Config");
    jmethodID valueOfBitmapConfigFunction = env->GetStaticMethodID(bitmapConfigClass, "valueOf", "(Ljava/lang/String;)Landroid/graphics/Bitmap$Config;");
    jobject bitmapConfig = env->CallStaticObjectMethod(bitmapConfigClass, valueOfBitmapConfigFunction, configName);
    jobject newBitmap = env->CallStaticObjectMethod(bitmapCls, createBitmapFunction, info.height, info.width, bitmapConfig);
    //
    // putting the pixels into the new bitmap:
    //
    if ((ret = AndroidBitmap_lockPixels(env, newBitmap, &bitmapPixels)) < 0)
    {
        LOGE("AndroidBitmap_lockPixels() failed ! error=%d", ret);
        return NULL;
    }
    uint32_t* newBitmapPixels = (uint32_t*) bitmapPixels;
    int whereToPut = 0;
 
    ////
    for (col = 0; col < mWidth; col++) {
 
        for (row = 0; row < mHeight; row++) {
            blue = 0;
            green = 0;
            red = 0;
            void *pixel = NULL;
            for (m = 0; m < 9; m++) {
                for (n = 0; n < 9; n++) {
                    y = (row + m - 4) % mHeight;
                    x = (col + n - 4) % mWidth;
                    if ((row + m - 4) < 0 || y >= mHeight || (col + n - 4) < 0 || x >= mWidth) continue;
                    LOGE("handleBitmapForSinglePixel %d", pixelArr[0])
                    // Move pixel pointer
                    //pixel = pixelArr + y * mWidth + x;
 
                    //Get the value corresponding to the pointer according to the ABGR storage sequence
                    //uint32_t v = *((uint32_t *) pixel);
                    uint32_t v = tempPixels[info.width * y + x];
 
                    //
                    a = RGB8888_A(v);
                    r = RGB8888_R(v);
                    g = RGB8888_G(v);
                    b = RGB8888_B(v);
 
 
                    red += r * mask[m][n];
                    green += g * mask[m][n];
                    blue += b * mask[m][n];
                }
            }
            *((uint32_t *) pixel) = MAKE_ABGR(a, red, green, blue);
            newBitmapPixels[whereToPut++] = pixel;
        }
    }
 
    AndroidBitmap_unlockPixels(env, newBitmap);
    //
    // freeing the native memory used to store the pixels
    //
    free(tempPixels);
    return newBitmap;
}

  ver 2

//
// Created by kyun1016 on 2020-11-18.
//
#include<jni.h>
#include<stdio.h>
#include<stdlib.h>
#include <sys/time.h>
#include <android/log.h>
#include <android/bitmap.h>
 
#include <CL/opencl.h>
#include <assert.h>
 
#define CL_FILE "/data/local/tmp/blur.cl"
 
#define LOG_TAG "DEBUG"
#define LOGD(...) __android_log_print(ANDROID_LOG_DEBUG,LOG_TAG,__VA_ARGS__)
#define LOGE(...) __android_log_print(ANDROID_LOG_ERROR,LOG_TAG,__VA_ARGS__)
 
#define RGB8888_A(p) ((p & (0xff<<24))      >> 24 )
#define RGB8888_B(p) ((p & (0xff << 16)) >> 16 )
#define RGB8888_G(p) ((p & (0xff << 8))  >> 8 )
#define RGB8888_R(p) (p & (0xff) )
 
#define CHECK_CL(err) {\
    cl_int er = (err);\
    if(er<0 && er > -64){\
        LOGE("%d line, OpenCL Error:%d\n",__LINE__,er);\
    }\
}
 
JNIEXPORT jobject JNICALL
Java_com_example_week12_1opencl_1with_1cpp_MainActivity_GaussianBlurBitmap
        (JNIEnv *env, jclass class , jobject bitmap)
{
    //getting bitmap info:
    LOGD("reading bitmap info...");
    AndroidBitmapInfo info;
    int ret;
    if ((ret = AndroidBitmap_getInfo(env, bitmap, &info)) < 0) {
        LOGE("AndroidBitmap_getInfo() failed ! error=%d", ret);
        return NULL;
    }
    LOGD("width:%d height:%d stride:%d", info.width, info.height, info.stride);
    if (info.format != ANDROID_BITMAP_FORMAT_RGBA_8888) {
        LOGE("Bitmap format is not RGBA_8888!");
        return NULL;
    }
 
 
    //read pixels of bitmap into native memory :
    LOGD("reading bitmap pixels...");
    void* bitmapPixels;
    if ((ret = AndroidBitmap_lockPixels(env, bitmap, &bitmapPixels)) < 0) {
        LOGE("AndroidBitmap_lockPixels() failed ! error=%d", ret);
        return NULL;
    }
    uint32_t* src = (uint32_t*) bitmapPixels;
    uint32_t* tempPixels = (uint32_t*)malloc(info.height * info.width*4);
    int pixelsCount = info.height * info.width;
    memcpy(tempPixels, src, sizeof(uint32_t) * pixelsCount);
 
 
    // write image processing code here!!
    // input is tempPixels
    // output is bitmapPixels
 
    int a,r, g, b;
 
    //
    float red = 0, green = 0, blue = 0;
    int row = 0, col = 0, cnt=0;
    int m, n, k,x,y;
 
    float mask[9][9] = {
            {0.011237, 0.011637, 0.011931, 0.012111, 0.012172, 0.012111, 0.011931, 0.011637, 0.011237},
            {0.011637, 0.012051, 0.012356, 0.012542, 0.012605, 0.012542, 0.012356, 0.012051, 0.011637},
            {0.011931, 0.012356, 0.012668, 0.012860, 0.012924, 0.012860, 0.012668, 0.012356, 0.011931},
            {0.012111, 0.012542, 0.012860, 0.013054, 0.013119, 0.013054, 0.012860, 0.012542, 0.012111},
            {0.012172, 0.012605, 0.012924, 0.013119, 0.013185, 0.013119, 0.012924, 0.012605, 0.012172},
            {0.012111, 0.012542, 0.012860, 0.013054, 0.013119, 0.013054, 0.012860, 0.012542, 0.012111},
            {0.011931, 0.012356, 0.012668, 0.012860, 0.012924, 0.012860, 0.012668, 0.012356, 0.011931},
            {0.011637, 0.012051, 0.012356, 0.012542, 0.012605, 0.012542, 0.012356, 0.012051, 0.011637},
            {0.011237, 0.011637, 0.011931, 0.012111, 0.012172, 0.012111, 0.011931, 0.011637, 0.011237},
    };
 
    for (col = 0; col < info.width; col++) {
 
        for (row = 0; row < info.height; row++) {
            blue = 0;
            green = 0;
            red = 0;
 
            for (m = 0; m < 9; m++) {
                for (n = 0; n < 9; n++) {
                    y = (row + m - 4);
                    x = (col + n - 4);
                    if ((row + m - 4) < 0 || y >= info.height || (col + n - 4) < 0 || x >= info.width) continue;
                    uint32_t pixel = tempPixels[info.width * y + x];
 
                    //
                    a = RGB8888_A(pixel);
                    r = RGB8888_R(pixel);
                    g = RGB8888_G(pixel);
                    b = RGB8888_B(pixel);
 
 
                    red += r * mask[m][n];
                    green += g * mask[m][n];
                    blue += b * mask[m][n];
                    //LOGD("r : %f   g : %f   b : %f", red, green, blue);
                }
            }
            r = (int) red;
            g = (int) green;
            b = (int) blue;
            uint32_t v = (a << 24) + (b << 16) + (g << 8) + (r);
            src[info.width * row + col] = v;
            cnt++;
        }
    }
 
    AndroidBitmap_unlockPixels(env, bitmap);
    //
    // free the native memory used to store the pixels
    //
    free(tempPixels);
    return bitmap;
}
 
 
 
 
 
JNIEXPORT jobject JNICALL
Java_com_example_week12_1opencl_1with_1cpp_MainActivity_GaussianBlurGPU
        (JNIEnv *env, jclass class, jobject bitmap )
{
    //getting bitmap info:
    LOGD("reading bitmap info...");
    AndroidBitmapInfo info;
    int ret;
    if ((ret = AndroidBitmap_getInfo(env, bitmap, &info)) < 0) {
        LOGE("AndroidBitmap_getInfo() failed ! error=%d", ret);
        return NULL;
    }
    LOGD("width:%d height:%d stride:%d", info.width, info.height, info.stride);
    if (info.format != ANDROID_BITMAP_FORMAT_RGBA_8888) {
        LOGE("Bitmap format is not RGBA_8888!");
        return NULL;
    }
 
 
    //read pixels of bitmap into native memory :
    LOGD("reading bitmap pixels...");
    void* bitmapPixels;
    if ((ret = AndroidBitmap_lockPixels(env, bitmap, &bitmapPixels)) < 0) {
        LOGE("AndroidBitmap_lockPixels() failed ! error=%d", ret);
        return NULL;
    }
    uint32_t* src = (uint32_t*) bitmapPixels;
    uint32_t* tempPixels = (uint32_t*)malloc(info.height * info.width*4);
    int pixelsCount = info.height * info.width;
    memcpy(tempPixels, src, sizeof(uint32_t) * pixelsCount);
 
    ////GPU
    LOGD("GPU Start");
    FILE* file_handle;
    char* kernel_file_buffer, * file_log;;
    size_t kernel_file_size, log_size;
 
    unsigned char* cl_file_name = CL_FILE;
    unsigned char* kernel_name = "kernel_blur";
 
    //Device input buffers
    cl_mem d_src;
    //Device output buffer
    cl_mem d_dst;
 
    cl_platform_id clPlatform;        //OpenCL platform
    cl_device_id device_id;            //device ID
    cl_context context;                //context
    cl_command_queue queue;            //command queue
    cl_program program;                //program
    cl_kernel kernel;                //kernel
    LOGD("cl_file_open");
    file_handle = fopen(cl_file_name, "r");
    if (file_handle == NULL) {
        printf("Couldn't find the file");
        exit(1);
    }
 
    //read kernel file
    fseek(file_handle, 0, SEEK_END);
    kernel_file_size = ftell(file_handle);
    rewind(file_handle);
    kernel_file_buffer = (char*)malloc(kernel_file_size + 1);
    kernel_file_buffer[kernel_file_size] = '\0';
    fread(kernel_file_buffer, sizeof(char), kernel_file_size, file_handle);
    fclose(file_handle);
    LOGD("%s",kernel_file_buffer);
    LOGD("file_buffer_read");
    // Initialize vectors on host
    int i;
 
    size_t globalSize, localSize, grid;
    cl_int err;
 
    // Number of work items in each local work group
    localSize = 64;
    int n_pix = info.width * info.height;
 
    //Number of total work items - localSize must be devisor
    grid = (n_pix % localSize) ? (n_pix / localSize) + 1 : n_pix / localSize;
    globalSize = grid * localSize;
 
    LOGD("calc grid and globalSize");
    //openCL 기반 실행
 
    //Bind to platform
    LOGD("error check");
    CHECK_CL(clGetPlatformIDs(1, &clPlatform, NULL));
    LOGD("error end check");
    //Get ID for the device
    CHECK_CL(clGetDeviceIDs(clPlatform, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL));
    //Create a context
    context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
    CHECK_CL(err);
    //Create a command queue
    queue = clCreateCommandQueue(context, device_id, 0, &err);
    CHECK_CL(err);
    //Create the compute program from the source buffer
    program = clCreateProgramWithSource(context, 1, (const char**)&kernel_file_buffer, &kernel_file_size, &err);
    CHECK_CL(err);
    //Build the program executable
    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    LOGD("error 22 check");
    CHECK_CL(err);
    if (err != CL_SUCCESS) {
        //LOGD("%s", err);
        size_t len;
        char buffer[4096];
        LOGD("Error: Failed to build program executable!");
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer),
                              buffer, &len);
 
        LOGD("%s", buffer);
        //exit(1);
    }
    LOGD("error 323 check");
 
    //Create the compute kernel in the program we wish to run
    kernel = clCreateKernel(program, kernel_name, &err);
    CHECK_CL(err);
 
 
 
 
    //////openCL 커널 수행
    //Create the input and output arrays in device memory for our calculation
    d_src = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(uint32_t)*info.width*info.height, NULL, NULL);
    d_dst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(uint32_t)*info.width*info.height, NULL, NULL);
 
    //Write our data set into the input array in device memory
    CHECK_CL(clEnqueueWriteBuffer(queue, d_src, CL_TRUE, 0, sizeof(uint32_t)*info.width*info.height, tempPixels, 0, NULL, NULL));
 
    //Set the arguments to our compute kernel
    CHECK_CL(clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_src));
    CHECK_CL(clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_dst));
    CHECK_CL(clSetKernelArg(kernel, 2, sizeof(uint32_t), &info.width));
    CHECK_CL(clSetKernelArg(kernel, 3, sizeof(uint32_t), &info.height));
 
 
    //Execute the kernel over the entire range of the data set
    CHECK_CL(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize, 0, NULL, NULL));
    //Wait for the command queue to get serviced before reading back results
    CHECK_CL(clFinish(queue));
    //read the results form the device
    CHECK_CL(clEnqueueReadBuffer(queue, d_dst, CL_TRUE, 0, sizeof(uint32_t)*info.width*info.height, src, 0, NULL, NULL));
 
 
    // release OpenCL resources
    CHECK_CL(clReleaseMemObject(d_src));
    CHECK_CL(clReleaseMemObject(d_dst));
    CHECK_CL(clReleaseProgram(program));
    CHECK_CL(clReleaseKernel(kernel));
    CHECK_CL(clReleaseCommandQueue(queue));
    CHECK_CL(clReleaseContext(context));
 
    AndroidBitmap_unlockPixels(env, bitmap);
    //
    // free the native memory used to store the pixels
    //
    free(tempPixels);
    return bitmap;
 
    return 0;
}
 
 

- MainActivity.java

package com.example.week12_opencl_with_cpp;
 
import androidx.appcompat.app.AppCompatActivity;
 
import android.graphics.Bitmap;
import android.graphics.BitmapFactory;
import android.os.Bundle;
import android.view.View;
import android.widget.Button;
import android.widget.ImageView;
import android.widget.TextView;
 
public class MainActivity extends AppCompatActivity {
 
    // Used to load the 'native-lib' library on application startup.
    static {
        System.loadLibrary("OpenCLDriver");
    }
    //blur CPU
    public native Bitmap GaussianBlurBitmap(Bitmap bitmap);
    //blur GPU
    public native Bitmap GaussianBlurGPU(Bitmap bitmap);
 
    ImageView imgV;
    Bitmap buf_bitmap;
 
    @Override
    protected void onCreate(Bundle savedInstanceState) {
        super.onCreate(savedInstanceState);
        setContentView(R.layout.activity_main);
 
        final TextView tv = (TextView) findViewById(R.id.txt);
        imgV = (ImageView) findViewById(R.id.imageView);
        Button btn = (Button) findViewById(R.id.button);
        Button btn2 = (Button) findViewById(R.id.button2);
        Button btn3 = (Button) findViewById(R.id.button3);
 
        BitmapFactory.Options options = new BitmapFactory.Options();
        options.inPreferredConfig = Bitmap.Config.ARGB_8888;
 
        buf_bitmap=BitmapFactory.decodeFile("/data/local/tmp/lena.bmp", options);
 
        tv.setText("Image Processing App");
 
        imgV.setImageBitmap(buf_bitmap);
 
        btn.setOnClickListener(new View.OnClickListener() {
            @Override
            public void onClick(View view) {
                //When this button is clicked,
                //Process the Bitmap class using native function
                float start = (float)System.nanoTime()/1000000;
                buf_bitmap = GaussianBlurBitmap(buf_bitmap);
                //then set the result Bitmap image to the ImageView
                float end = (float)System.nanoTime()/1000000;
                float timesub = end-start;
                imgV.setImageBitmap(buf_bitmap);
                //and print the execution time on TextView
                String ts = Float.toString(timesub);
                tv.setText("Execution Time: " + ts + "ms");
 
            }
        });
        btn2.setOnClickListener(new View.OnClickListener() {
            @Override
            public void onClick(View view) {
                //When this button is clicked,
                //Process the Bitmap class using native function
                float start = (float)System.nanoTime()/1000000;
                buf_bitmap = GaussianBlurGPU(buf_bitmap);
                //then set the result Bitmap image to the ImageView
                float end = (float)System.nanoTime()/1000000;
                float timesub = end-start;
                imgV.setImageBitmap(buf_bitmap);
                //and print the execution time on TextView
 
                String ts = Float.toString(timesub);
                tv.setText("Execution Time: " + ts + "ms");
            }
        });
        btn3.setOnClickListener(new View.OnClickListener() {
            @Override
            public void onClick(View view) {
                BitmapFactory.Options options = new BitmapFactory.Options();
                options.inPreferredConfig = Bitmap.Config.ARGB_8888;
                buf_bitmap =
                        BitmapFactory.decodeFile("/data/local/tmp/lena.bmp",
                                options);
                imgV.setImageBitmap(buf_bitmap);
                tv.setText("Orignal image");
            }
        });
    }
 
}

- MainActivity.xml

<?xml version="1.0" encoding="utf-8"?>
<LinearLayout xmlns:android="http://schemas.android.com/apk/res/android"
    android:orientation="vertical"
    android:layout_width="match_parent"
    android:layout_height="match_parent">
 
    <ImageView
        android:id="@+id/imageView"
        android:layout_width="match_parent"
        android:layout_height="400dp" />
    <LinearLayout
        android:layout_width="match_parent"
        android:layout_height="wrap_content"
        android:gravity="center">
 
        <Button
            android:id="@+id/button"
            android:layout_width="wrap_content"
            android:layout_height="wrap_content"
            android:layout_gravity="center_vertical"
            android:text="CPU BLUR"/>
        <Button
            android:id="@+id/button2"
            android:layout_width="wrap_content"
            android:layout_height="wrap_content"
            android:layout_gravity="center_vertical"
            android:layout_marginLeft="20dp"
            android:text="GPU BLUR"/>
        <Button
            android:id="@+id/button3"
            android:layout_width="wrap_content"
            android:layout_height="wrap_content"
            android:layout_gravity="center_vertical"
            android:layout_marginLeft="20dp"
            android:text="ORIGINAL"/>
 
    </LinearLayout>
    <TextView
        android:id="@+id/txt"
        android:layout_width="wrap_content"
        android:layout_height="wrap_content"
        android:layout_marginLeft="20dp"
        android:textSize="20dp"
        android:text="Kyun123"/>
    <TextView
        android:layout_width="wrap_content"
        android:layout_height="wrap_content"
        android:layout_marginLeft="20dp"
        android:textSize="20dp"
        android:text="=== Kyun 2016440047 ==="/>
 
 
</LinearLayout>

- OpenCL

 1) Gaussian Blur

//OpenCL kernel. Each work item takes care of one element of c
 
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
 
#define RGB8888_A(p) ((p & (0xff<<24))      >> 24 )
#define RGB8888_B(p) ((p & (0xff << 16)) >> 16 )
#define RGB8888_G(p) ((p & (0xff << 8))  >> 8 )
#define RGB8888_R(p) (p & (0xff) )
 
__kernel void kernel_blur(__global int *src, __global int *dst, const int width, const int height){
    //int id = get_global_id(0);
    
    int row = get_global_id(0)/width;
    int col = get_global_id(0)%width;
 
    float mask[9][9] = {
        {0.011237, 0.011637, 0.011931, 0.012111, 0.012172, 0.012111, 0.011931, 0.011637, 0.011237},
        {0.011637, 0.012051, 0.012356, 0.012542, 0.012605, 0.012542, 0.012356, 0.012051, 0.011637},
        {0.011931, 0.012356, 0.012668, 0.012860, 0.012924, 0.012860, 0.012668, 0.012356, 0.011931},
        {0.012111, 0.012542, 0.012860, 0.013054, 0.013119, 0.013054, 0.012860, 0.012542, 0.012111},
        {0.012172, 0.012605, 0.012924, 0.013119, 0.013185, 0.013119, 0.012924, 0.012605, 0.012172},
        {0.012111, 0.012542, 0.012860, 0.013054, 0.013119, 0.013054, 0.012860, 0.012542, 0.012111},
        {0.011931, 0.012356, 0.012668, 0.012860, 0.012924, 0.012860, 0.012668, 0.012356, 0.011931},
        {0.011637, 0.012051, 0.012356, 0.012542, 0.012605, 0.012542, 0.012356, 0.012051, 0.011637},
        {0.011237, 0.011637, 0.011931, 0.012111, 0.012172, 0.012111, 0.011931, 0.011637, 0.011237},
    };
 
    int a,r, g, b;
 
    float red = 0, green = 0, blue = 0;
    int m, n,x,y;
 
    for (m = 0; m < 9; m++) {
        for (n = 0; n < 9; n++) {
            y = (row + m - 4);
            x = (col + n - 4);
            if ((row + m - 4) < 0 || y >= height || (col + n - 4) < 0 || x >= width) continue;
            int pixel = src[width * y + x];
 
            //
            a = RGB8888_A(pixel);
            r = RGB8888_R(pixel);
            g = RGB8888_G(pixel);
            b = RGB8888_B(pixel);
 
 
            red += r * mask[m][n];
            green += g * mask[m][n];
            blue += b * mask[m][n];
        }
    }
    r = (int) red;
    g = (int) green;
    b = (int) blue;
 
    int v = (a << 24) + (b << 16) + (g << 8) + (r);
    dst[width * row + col] = v;
}

2) GrayScale

//OpenCL kernel. Each work item takes care of one element of c
 
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
 
#define RGB8888_A(p) ((p & (0xff<<24))      >> 24 )
#define RGB8888_B(p) ((p & (0xff << 16)) >> 16 )
#define RGB8888_G(p) ((p & (0xff << 8))  >> 8 )
#define RGB8888_R(p) (p & (0xff) )
 
__kernel void kernel_gray(__global int *src, __global int *dst, const int width, const int height){
    //int id = get_global_id(0);
    
    int row = get_global_id(0)/width;
    int col = get_global_id(0)%width;
 
    int a,r, g, b;
    int gray;
 
    float red = 0, green = 0, blue = 0;
    int pixel = src[width * row + col];
    a = RGB8888_A(pixel);
    r = RGB8888_R(pixel);
    g = RGB8888_G(pixel);
    b = RGB8888_B(pixel);
 
 
    red = r * 0.2126;
    green = g * 0.7152;
    blue = b * 0.0722;
    gray = red+green+blue;
 
    int v = (a << 24) + (gray << 16) + (gray << 8) + (gray);
    dst[width * row + col] = v;
}

3) Rotate

//OpenCL kernel. Each work item takes care of one element of c
 
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
 
#define RGB8888_A(p) ((p & (0xff<<24))      >> 24 )
#define RGB8888_B(p) ((p & (0xff << 16)) >> 16 )
#define RGB8888_G(p) ((p & (0xff << 8))  >> 8 )
#define RGB8888_R(p) (p & (0xff) )
 
 
__kernel void kernel_rot(__global int *src, __global int *dst, const int width, const int height){
    //int id = get_global_id(0);
    
    int row = get_global_id(0)/width;
    int col = get_global_id(0)%width;
 
    dst[width * col + row] = src[width * row + col];
}

3. 참고

https://lsit81.tistory.com/entry/JNI-%EA%B8%B0%EC%B4%88

 

https://realapril.tistory.com/35

 

https://webnautes.tistory.com/1090

 

https://stackoverflow.com/questions/14398670/rotating-a-bitmap-using-jni-ndk

https://jasonjason.tistory.com/15

- libGLES_mali.so

https://answer-id.com/ko/9005352

질문이나 지적 있으시면 댓글로 남겨주세요~

도움 되셨으면 하트 꾹!