Android OpenCL测试程序，使用dlopen动态加载libOpenCL.so库

最新推荐文章于 2024-08-02 11:05:37 发布

wujushan

最新推荐文章于 2024-08-02 11:05:37 发布

阅读量1.2w

点赞数 1

分类专栏： Opencv4Android 文章标签： OpenCV OpenCL Opencv4Android Android Opencl

本文链接：https://blog.csdn.net/wjskeepmaking/article/details/70080315

版权

Opencv4Android 专栏收录该内容

6 篇文章 1 订阅

订阅专栏

1、环境配置：

Android NDK

Android studio

2、下载头文件：

大家可以通过这个链接下载头文件：

http://download.csdn.net/detail/wjskeepmaking/9811058

3、声明native函数并生成头文件：

public class OpenclTest {


    public String testOpencl(){
        return testopencl();
    }
    public String getplatformName(){
        return getPlatformName();
    }
    public String getdeviceName(){
        return getDeviceName();
    }

    public native String  testopencl();
    public native String getPlatformName();
    public native String getDeviceName();
}

进入package的同级目录，利用javah命令生成头文件：

4、编写ocl.cpp文件：

//
// Created by wujs on 2017/4/11.
//
#include "CL/cl.h"
#include<malloc.h>
#include<stdio.h>
#include<stdlib.h>
#include <dlfcn.h>
#include "topencl.h"
#include"com_pax_imagesobelfilter_OpenclTest.h"
#include <android/log.h>
#define LOG_TAG "test"
#define LOGD(...) ((void)__android_log_print(ANDROID_LOG_DEBUG, LOG_TAG, __VA_ARGS__))
#define LEN(arr) sizeof(arr) / sizeof(arr[0])
#define N 1024
#define NUM_THREAD 128

cl_uint num_device;
cl_uint num_platform;
cl_platform_id *platform;
cl_device_id *devices;
cl_int err;
cl_context context;
cl_command_queue cmdQueue;
cl_mem buffer,sum_buffer;
cl_program program ;
cl_kernel kernel;
const char* src[] = {
        "  __kernel void redution(  \n"
                "  __global int *data,     \n"
                "  __global int *output,   \n"
                "  __local int *data_local   \n"
                "  )  \n"
                " {   \n"
                "  int gid=get_group_id(0);   \n"
                "  int tid=get_global_id(0);    \n"
                "  int size=get_local_size(0);   \n"
                "  int id=get_local_id(0);     \n"
                "  data_local[id]=data[tid];   \n"
                "  barrier(CLK_LOCAL_MEM_FENCE);   \n"
                "  for(int i=size/2;i>0;i>>=1){    \n"
                "      if(id<i){   \n"
                "          data_local[id]+=data_local[id+i];   \n"
                "      }   \n"
                "      barrier(CLK_LOCAL_MEM_FENCE);   \n"
                "  }    \n"
                "  if(id==0){    \n"
                "      output[gid]=data_local[0];   \n"
                "  }    \n"
                " }   \n"

};
int num_block;
int  test()
{
    int* in,*out;
    num_block=N/NUM_THREAD;
    in=(int*)malloc(sizeof(int)*N);
    out=(int*)malloc(sizeof(int)*num_block);
    for(int i=0;i<N;i++){
        in[i]=1;
    }
    initFns(); //load library
    Init_OpenCL();
    Context_cmd();
    Create_Buffer(in);
    Create_program();
    Set_arg();
    Execution();
    CopyOutResult(out);
    int sum=0;
    for(int i=0;i<num_block;i++){
        sum+=out[i];
    }
    return sum;
}

JNIEXPORT jstring JNICALL Java_com_pax_imagesobelfilter_OpenclTest_testopencl(JNIEnv * env, jobject thisobject)
{
    char result[10];
    LOGD("test_opencl start to test");
    sprintf(result,"%d\n",test());
    return env->NewStringUTF(result);
}
JNIEXPORT jstring JNICALL Java_com_pax_imagesobelfilter_OpenclTest_getPlatformName(JNIEnv *env , jobject thisobject)
{
    char buffer[1024];
    LOGD("getPlatformName start");
    clGetPlatformInfo(platform[0],CL_PLATFORM_NAME,sizeof(buffer),buffer,NULL);
    return env->NewStringUTF(buffer);
}
JNIEXPORT jstring JNICALL Java_com_pax_imagesobelfilter_OpenclTest_getDeviceName(JNIEnv *env , jobject thisobject)
{

    char buffer[1024];
    LOGD("getDeviceName start");
    clGetDeviceInfo(devices[0],CL_DEVICE_NAME,sizeof(buffer),buffer,NULL);
    return env->NewStringUTF(buffer);
}

void Init_OpenCL()
{
    LOGD("start init OpenCL");
    size_t nameLen1;
    char platformName[1024];

    err = clGetPlatformIDs(0, 0, &num_platform);
    LOGD("err=clGetPlatformIDs");
    platform=(cl_platform_id*)malloc(sizeof(cl_platform_id)*num_platform);
    err = clGetPlatformIDs(num_platform, platform, NULL);

    err=clGetDeviceIDs(platform[0],CL_DEVICE_TYPE_GPU,0,NULL,&num_device);
    devices=(cl_device_id*)malloc(sizeof(cl_device_id)*num_device);
    err=clGetDeviceIDs(platform[0],CL_DEVICE_TYPE_GPU,num_device,devices,NULL);

}

void Context_cmd()
{
    context=clCreateContext(NULL,num_device,devices,NULL,NULL,&err);
    cmdQueue=clCreateCommandQueue(context,devices[0],0,&err);
}

void Create_Buffer(int *data)
{

    buffer=clCreateBuffer(context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,sizeof(int)*N,data,&err);
    sum_buffer=clCreateBuffer(context,CL_MEM_WRITE_ONLY,sizeof(int)*num_block,0,&err);
}

void Create_program()
{
    program=clCreateProgramWithSource(context, LEN(src), src, NULL, NULL);
    err=clBuildProgram(program,num_device,devices,NULL,NULL,NULL);
    kernel = clCreateKernel(program, "redution", NULL);
}

void Set_arg()
{
    err=clSetKernelArg(kernel,0,sizeof(cl_mem),&buffer);
    err=clSetKernelArg(kernel,1,sizeof(cl_mem),&sum_buffer);
    err=clSetKernelArg(kernel,2,sizeof(int)*NUM_THREAD,NULL);
}

void Execution()
{
    const size_t globalWorkSize[1]={N};
    const size_t localWorkSize[1]={NUM_THREAD};
    err=clEnqueueNDRangeKernel(cmdQueue,kernel,1,NULL,globalWorkSize,localWorkSize,0,NULL,NULL);
    clFinish(cmdQueue);
}

void CopyOutResult(int*out)
{
    err=clEnqueueReadBuffer(cmdQueue,sum_buffer,CL_TRUE,0,sizeof(int)*num_block,out,0,NULL,NULL);
}

5、创建topencl.h文件和topencl.c文件；

利用dlopen,实现动态加载libOpenCL.so库，由于代码太多，这里就不全部贴出来了，后面会有链接，大家可以去下载，根据自己的实际情况修改，需要修改的地方不会很多。

Note：下面代码是从系统中去加载动态库，这里有个问题，如果大家使用的机器是Android N以下，加载不会出错，但是，如果是Android N或以上的，加载会出现问题，因为从Android N开始，限制了程序对系统私有库（除了Android NDK中提供的那些库）的加载了。针对这种情况，现在的解决方法就是：把需要加载和依赖的库全部导入APP中，直接从APP中去加载，大家可以参阅这篇文章：解决Android N加载系统私有库的错误。

下面是topencl.c的库加载代码：

#define LIB_OPENCL "/system/vendor/lib/libOpenCL.so"
#define LIB_GLES_MALI "/system/vendor/lib/egl/libGLES_adreno.so"
#define LIB_LLVM   "/system/lib/libllvm-a3xx.so"

static void *getCLHandle(){

    LOGD("get_handle");
    void *res = NULL;
    char* so_name="Unknown Shared library for OpenCL";
    res = dlopen(LIB_OPENCL,RTLD_LAZY);
    if(res != NULL){
        so_name = LIB_OPENCL;
    }else{
        res = dlopen(LIB_GLES_MALI,RTLD_LAZY);
    }
    if(res != NULL){
        so_name = LIB_GLES_MALI;
    }else{
        res = dlopen(LIB_LLVM,RTLD_LAZY);
    }
    /*  if(res==NULL){
      res = dlopen("/home/rahul/stream/sdk2.7/lib/x86_64/libOpenCL.so",RTLD_LAZY);
    }*/
    if(res != NULL) {
        so_name=LIB_LLVM;
    } else {
        LOGD("Could not open library :(\n");
    }
    LOGD("Loaded library name:%s\n",so_name);
    return res;
}

6、编写Android.mk文件和Application.mk文件：

android.mk文件：

要根据自己的Opencv4Android的路径去修改OPENCV_ANDROID_SDK

LOCAL_PATH := $(call my-dir)

include $(CLEAR_VARS)
OPENCV_ANDROID_SDK := F:\Android\opencv-3.2.0-android-sdk\OpenCV-android-sdk
ifdef OPENCV_ANDROID_SDK
  ifneq ("","$(wildcard $(OPENCV_ANDROID_SDK)/OpenCV.mk)")
    include ${OPENCV_ANDROID_SDK}/OpenCV.mk
  else
    include ${OPENCV_ANDROID_SDK}/sdk/native/jni/OpenCV.mk
  endif
else
  include ../../sdk/native/jni/OpenCV.mk
endif

LOCAL_MODULE    := ocl
LOCAL_CPPFLAGS	+= -I$(LOCAL_PATH)
LOCAL_SRC_FILES := ocl.cpp topencl.c
LOCAL_LDLIBS += -llog -ldl
include $(BUILD_SHARED_LIBRARY)

Application.mk文件：

APP_CPPFLAGS := -frtti -fexceptions
#APP_STL := stlport_static
APP_STL := gnustl_static
APP_ABI := armeabi-v7a arm64-v8a
APP_PLATFORM := android-8

7、编译：

进入 jni 的同级目录，利用ndk-build进行编译：

、

8、程序运行结果：

以上例程是参考这篇博文写的：http://blog.csdn.net/wcj0626/article/details/26272019#comments。但这个链接中使用静态库进行编译的。所以，我在这里写一个通过动态加载手机中的OpenCL库的文章，供大家参考使用。我的另一篇文章，使用opencv T-API来获取GPU信息： Opencv4Android的OpenCL的测试，使用Opencv的ocl封装库

代码链接： CSDN

OpenCV manager

因为Opencv4Android sdk 太大了，所以，我从代码中剔除之后再上传的，如果大家下载学习，要把Opencv4Android SDK 导入和安装OpenCV manager.apk（上面提供了连链接下载，除了X86和X86_64）之后才能正确运行。

如果大家有什么疑问，也请在评论中提出，能回答尽量回答，大家互相学习。