前言

Metal的官方文档中有一个利用GPU进行计算的例子：Performing Calculations on a GPU。本文主要探讨的是如何在Android上实现类似的效果。

使用OpenCL获得GPU的能力

之所以选用OpenCL，是因为它的api比较符合通用计算的场景。虽然它不仅适前端性能优化合approveGPU，还适合其他单元。但本文主要会利用OpenCL来进行GPU运算。

GPU 优化技术-OpenCL 介绍

由于NDK并没有提供关于O测试抑郁症penCL的动态库，所以在Android中使用OpenCL需要弄一些骚操作。

从手机上导出动态库并在工程中使用

以笔者使用的测测试用例试机Oppo R15为例，OpenCL环境可在设备的/system/vendor/lib/libOpenCL.so获取。此时可通过adb pull公司让员工下班发手机电量截图将so包导出，放入工程中使用。

adb pull /system/vendor/lib/libOpenCL.so

还有一个比较重要的是CL的头文枸杞件，这个可以在官方的github上下载：

c的头文件：KhronosGr测试你适合学心理学吗oup/OpenCL-Headers
c++的公积金封装：KhronosGroup/OpenCL-CLHPP

需要注意的是：

由于OpenCL的版本问题，上述头文件定义的某些逻辑，对github汤姆应到您在设备上导出github官网登陆入口的so包不一定试用。需要按OpenCL真实的版本进行代码编写。
由于so包是在单一设备导出的，不排除在其他设备有适配问题，且一般高通平台Adreno gpu命测试用例名为libOpenCL.soGo，但其他平台则不approve然测试你适合学心理学吗。

此部分可参考：高通平台下安卓opencl小例子

使用dlopen

笔者希appear望能够尽可能多的在不同设备上运行，所以想尽量使用在设备上的环境。

参考csarron/cl_stub: Ease OpenCL library loading可知，我们可以通过dlopgithub是什么en来获取OpenCL环境以及特定方法的句柄。简单分析一下该库的做法。

cl_stub

首先appointment还是需要将OpenCL的头文件导性能优化入到项目中：

c的头文件：KhronosGroup/OpenCL-Headers
c++的封装：KhronosGroup/OpenCL-CLHPP

定义一个可能的Open品悟性能优化CL动态库路径数组：

static const char *default_so_paths[] = {
        "/system/lib64/libOpenCL.so",
        "/system/lib64/egl/libGLES_mali.so",
        "/system/lib64/libPVROCL.so",
        "/system/vendor/lib64/libOpenCL.so",
        "/system/vendor/lib64/egl/libGLES_mali.so",
        "/system/vendor/lib64/libPVROCL.so",
        "libOpenCL.so"
};

在需要获取OpenCL环境测试你适合学心理学吗时，循环遍历该数组，寻找一个可以读取的路径。如果可以读取则被认为是可用性能优化于加载的属于该设备的Op测试英文enCL动态库路径。

static int access_file(const char *filename) {
    struct stat buffer;
    return (stat(filename, &buffer) == 0);
}
static int open_libopencl_so() {
    char *path = NULL
    ...
    if (!path) {
        for (i = 0; i < (sizeof(default_so_paths) / sizeof(char *)); i++) {
            if (access_file(default_so_paths[i])) {
                path = (char *) default_so_paths[i];
                break;
            }
        }

ps：该库的这里在找到path之后，测试你适合学心理学吗会有一个尝试将patappearanceh的路径替换成apk下路径的操作，目的是为了使用自己apk的so包，但该操作存在一些问题appear会导致dlopen返回NULL。建议将此工资超过5000怎么扣税段代码注释。

最后通过dlopen获取该库的实例，用于后面方法调用的使用。

so_handle = dlopen(path, RTLD_LAZY);

在该库中也有实现与OpenCL的api名字相同的函数，我们以获取平台品悟性能优化的api-clGetPlatformIDs为例：

cl_int
clGetPlatformIDs(cl_uint num_entries,
                 cl_platform_id *platforms,
                 cl_uint *num_platforms) {
    f_clGetPlatformIDs func;
    if (!so_handle)
        open_libopencl_so();
    func = (f_clGetPlatformIDs) dlsym(so_handle, "clGetPlatformIDs");
    if (func) {
        return func(num_entries, platforms, num_platforms);
    } else {
        return CL_INVALID_PLATFORM;
    }
}

这里通过上述事例出来的实apple例so_handle，利用dlsym获取到clGetPlatformIDs的句柄，继而调用它。至此就完成了一次调用OpenCL的api对接。

ps：f_clGetPlatformIDs为定义的一个与clGetPlatformIDs参数结构一致的函数别名

typedef cl_int (*f_clGetPlatformIDs)(cl_uint, cl_platform_id *, cl_uint *);

整体效果有点类似代理的味道，该做法类似在自己的库中实现一次OpenCL头文件定义的函数，利用dlopen/dlsym，达到从自己到OpenCL的函数调用效果。

而在外部调用时，则可前端性能优化以直接调用OpenCL头文件（CL/cl.h）定义的函数，或者C++封装类（CL/cl.hpp），最终都会走到像上述定义的函数中通过dlsym获取句柄。ps测试你的自卑程度：调用的逻辑会在后面的代码片段贴出。

扩展

类似品悟性能优化cl_stub的apple做法，在alibaba/MNN使用OpenCL能力的代码中也有体现。

在approveMNN/Open测试用例CLWrapper.cppapple有类似逻辑：

static const std::vector<std::string> gOpencl_library_paths = {
...
#elif defined(__ANDROID__)
    "libOpenCL.so",
    "libGLES_mali.so",
    "libmali.so",
#if defined(__aarch64__)
    // Qualcomm Adreno
    "/system/vendor/lib64/libOpenCL.so",
    "/system/lib64/libOpenCL.so",
    // Mali
    "/system/vendor/lib64/egl/libGLES_mali.so",
    "/system/lib64/egl/libGLES_mali.so",
#else
    // Qualcomm Adreno
    "/system/vendor/lib/libOpenCL.so", "/system/lib/libOpenCL.so",
    // Mali
    "/system/vendor/lib/egl/libGLES_mali.so", "/system/lib/egl/libGLES_mali.so",
    // other
    "/system/vendor/lib/libPVROCL.so", "/data/data/org.pocl.libs/files/lib/libpocl.so"
#endif
...
bool OpenCLSymbols::LoadLibraryFromPath(const std::string &library_path) {
#if defined(WIN32)
...
#else
    handle_ = dlopen(library_path.c_str(), RTLD_NOW | RTLD_LOCAL);
    if (handle_ == nullptr) {
        return false;
    }
#define MNN_LOAD_FUNCTION_PTR(func_name) func_name = reinterpret_cast<func_name##Func>(dlsym(handle_, #func_name)); 
    if(func_name == nullptr){ 
        mIsError = true; 
    }
#endif
    MNN_LOAD_FUNCTION_PTR(clGetPlatformIDs);
    MNN_LOAD_FUNCTION_PTR(clGetPlatformInfo);
    MNN_LOAD_FUNCTION_PTR(clBuildProgram);
    MNN_LOAD_FUNCTION_PTR(clEnqueueNDRangeKernel);
    MNN_LOAD_FUNCTION_PTR(clSetKernelArg);
    MNN_LOAD_FUNCTION_PTR(clReleaseKernel);
    MNN_LOAD_FUNCTION_PTR(clCreateProgramWithSource);
    MNN_LOAD_FUNCTION_PTR(clCreateBuffer);
    ...

ps：有兴趣的同学可以完整看看该目录下的封装：MNN/source/backend/opencl/core/runtime

ps：Android N之后，系统会限制获取私枸杞有库。所以需要注意您获取的OpenCL库是否存放在设备的工龄越长退休金越多吗系统私有库当中。Android 7.0 行为变更。

实现在GPU进行计算

以下代码笔者运用了OpenCL的C++封装，您也可以直接使github下载用其c代码的api效果是一致的。

clsl定义

在Metal的计算例子中，我们可以看到它定义的内核函数：

kernel void add_arrays(device const float* inA,
                       device const float* inB, 
                       device float* result, 
                       uint index [[thread_position_in_grid]]) { 
    result[index] = inA[index] + inB[index];
}

而在OpenCL中，我们也能定义一个clsl：

static const char* addClSl = R"clsl(
__kernel void add(__global const float *a, __global const float *b, __global float *c)
{
    int gid = get_global_id(0);
    c[gid] = a[gid] + b[gid];
}
)clsl";

当然，这个也能用.cl文件定义。

函数的定义为：a(inA)和b(inB)每项相加，得到c(resulappetitet)的每项。gid和Metal中的index定义都为该线程在网格中的位置。这里可将其表示为第几次调用该函数进行运算。

定龚俊义测试数据

static const int arrayLength = 1 << 24;
static const int bufferSize = arrayLength * sizeof(float);
float* data = (float *)malloc(bufferSize);
float* data2 = (float *)malloc(bufferSize);
float* result = (float *)malloc(bufferSize);
for (int i = 0; i < arrayLength; i++) {
    data[i] = (float)random()/(float)(RAND_MAX);
    data2[i] = (float)random()/(float)(RAND_MAX);
}

这里我龚俊们按APP照Metal的Demo，给数据的长度设置为1 <<测试你适合学心理学吗; 24，并为data和data2填龚俊入随机数。

创建环境

std::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.empty()) {
    LOGCATE("Platform size 0");
    return;
}
cl_context_properties properties[] =
        {CL_CONTEXT_PLATFORM, (cl_context_properties) (platforms[0])(), 0};
cl::Context context(CL_DEVICE_TYPE_GPU, properties);
std::vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
cl::CommandQueue queue(context, devices[0], 0, &err);
std::string kernelSource = addClSl;
cl::Program::Sources source(1, std::make_pair(kernelSource.c_str(),
                                              kernelSource.length() + 1));
cl::Program program(context, source);
const char *options = "-cl-fast-relaxed-math";
program.build(devices, options);
cl::Kernel kernel(program, "add", &err);

上述代码中做了几件事：

获取平台cl::Platform::get(&platforms);，一般数组只有一项，取第0位即可。
实例一个上下文，并指定需要使用GPU的type，CL_DEVICE_TYPE_GPU。ps：properties数组可根据不同平台设置特殊的常量。appearance这一块由于没有深入研究就不多细说了，测试你的自卑程度详细可参考M测试NN/OpenCLRuntime.cpp。
根据上下文获取设备context.getInfo<CL_CO工资超过5000怎么扣税NTEXT_DEVIappearanceCES>()，一般数组只有一项，取第0位即可。
根据上下文和设备创建一个CommandQueue，用于后续的函数执行。
利用上述c公司让员工下班发手机电量截图lsl，创建一个Program，并对clsl进行编译。
Kernel为clsl中函数的映射，譬如上述定义的add方法，则会被实例成KerneGitHubl。

数据写入

cl::Buffer aBuffer = cl::Buffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                bufferSize, (void *) &data[0], &err);
cl::Buffer bBuffer = cl::Buffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                bufferSize, (void *) &data2[0], &err);
cl::Buffer cBuffer = cl::Buffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                                bufferSize, (void *) &result[0], &err);
kernel.setArg(0, aBuffer);
kernel.setArg(1, bBuffer);
kernel.setArg(2, cBuffer);

上述逻辑将之前造的数据（applicationdata、data2、result）创建对应gpu能使用公司让员工下班发手机电量截图的Buffer。并将其设置到Kernel中作为函数的输入。对应的是clsl中的三个输入。

CL_MEM_COPY_HOST_PTR等flag的使用会关系到该内存在哪里创建的问题。这里的定义会在device中开辟一块内存（即所谓的显存），并将原有的内存（即电脑性能优化所谓的主存）赋值给他（data、data2、result）。这里会发生一次内存拷贝，属于耗时操作。

详细定义可参考：从零开始学习OpenCL开发（三）深入API

执行

size_t workgroupSize = 0;
err = kernel.getWorkGroupInfo(devices[0], CL_KERNEL_WORK_GROUP_SIZE, &workgroupSize);
err = queue.enqueueNDRangeKernel(kernel,
                                cl::NullRange,
                                cl::NDRange(arrayLength),
                                cl::NDRange(workgroupSize));
queue.finish();

将K前端性能优化ernel添加到CommandQueue中，由此可见gpu的执行对测试抑郁症的20道题于cpu是异步操作。最后可通过queue.finish();同步等待执行完成。

关于工作组和工作项的设置

工作组和工作项在Meapprovetal中也可定义为线google程组和线程，queappearue.enqueueNDRangeKernel需要设置全appearance部工作项数（global_work_size）和每个工作组的工作项数（local_work_size）。

一个工作组可容纳多少个工作Go项，这个是有最大值的，可通过Kernel获取。即上述的workgroupSize。
全部工作项数，因为这里我们的运算可以理解为循环相枸杞加1 << 24次，所以可以设置为1 << 24。

所以最后网格中有1 << 24个工作项github中文社区，（1 << 24）/ workgroupSize个宫颈癌工作组。因为这工商银行里的运算只是一维的，所以只需考虑一维即可。

此部分可参考：

OpenCl 对维度，工作项，工作组概念的理解

OpenCL与Metal API下如何合理地安排线程组大小

Metal框架详细解析APP（十五） —github汤姆— 计算处理之关于线程和线程组（三）

扩展阅读：

OpenCL优化：工作组大小性能优化

读取运算结approve果

queue.enqueueReadBuffer(cBuffer, CL_TRUE, 0, bufferSize, result);

从deivce中将运算结果拷贝回result，作为结果的测试抑郁程度的问卷输出读取。

至此，在Android上利用GPU进行运算的效果就基本实现了。Demo会贴在文末，以供参考工龄越长退休金越多吗。

扩展-使用OpenGL实现计算

OpenGL也是可以实现简单的数值运算的，只是由于它本身不是设计来支持简单的运算的。所以不能很好的适配。在OpenGL ES 3.0后可以利用Transform Feedback将在顶点着色器运算好的结果映射并取出。具体流程如下：

我们可以定义一个用于计算的顶点着色器和一个空的片元着色器：

static const GLchar* vertexShaderSrc = R"glsl(#version 300 es
layout(location = 0) in float inValue;
layout(location = 1) in float inValue2;
out float outValue;
void main()
{
    outValue = inValue + inValue2;
}
)glsl";
static const GLchar* fragmentShaderSrc = R"glsl(#version 300 es
void main() {
}
)glsl";

计算完成后可通过事先映射到Transform Feedback的outValue，读取出计算结果。这里可以使用EGL搭建OpeapproachnGL的环境。

详细可参考文章：

OpenGL – Transform Feedback

OpenGL ES 3.0使用Transform Feedback进行通用计算并读取结果

Android OpenGL ES 系列连载：（07）Transforappstorem Feedback ps：这里面的Demo有关于EGL的封装。

这部分的代码可参考Java实现：

Example of Android implementation of Transform Feedback in OpenGLES 3.0 and glMapBufferRange to read out data.

扩展-获取设备GPUGitHub基本信息

利用OpenCL可获取设备GPgithub中文社区U的基本信息，这里列举几个最基枸杞本的，譬如deviceName、deviceVersion。

std::vector<cl::Platform> platforms;
cl_int res = cl::Platform::get(&platforms);
if (platforms.empty() || res != CL_SUCCESS) {
    LOGCATE("Platform size 0");
    return;
}
const std::string platformName = platforms[0].getInfo<CL_PLATFORM_NAME>();
LOGCATD("cl_platform_name:%s", platformName.c_str());
std::vector<cl::Device> gpuDevices;
res = platforms[0].getDevices(CL_DEVICE_TYPE_GPU, &gpuDevices);
if (gpuDevices.empty() || res != CL_SUCCESS) {
    LOGCATE("Devices size 0");
    return;
}
const std::string deviceName = gpuDevices[0].getInfo<CL_DEVICE_NAME>();
LOGCATD("cl_device_name:%s", deviceName.c_str());
const std::string deviceVersion = gpuDevices[0].getInfo<CL_DEVICE_VERSION>();
LOGCATD("cl_device_version:%s", deviceVersion.c_str());
const std::string deviceVendor = gpuDevices[0].getInfo<CL_DEVICE_VENDOR>();
LOGCATD("cl_device_vendor:%s", deviceVendor.c_str());

Oppo R15上的输出结果：

D/gpu-compute: cl_platform_name:QUALCOMM Snapdragon(TM)
D/gpu-compute: cl_device_name:QUALCOMM Adreno(TM)
D/gpu-compute: cl_device_version:OpenCL 2.0 Adreno(TM) 512
D/gpu-compute: cl_device_vendor:QUALCOMM

这里推荐一个叫OpenCL-Z的应用，可以在Google Play上搜索得到，github官网安装后如果设备支持Open工商银行CL，可以在上面看到详细的设备信息以及不同平台对于它的扩展接口，以及库在系统的路径github下载。github直播平台永久回家

最后

本文主要介绍如何利用公司让员工下班发手机电量截图Ope性能优化的方法nCL在Android上实现Metal的计算逻辑，以及相关的环境适配与扩展。最后贴出上述的Demo，共同学习！ps：笔者使用的是ndk版本16.1.4479499，cmake版本3.6.4111459，可按需修改。

在Android上实现Metal的计算Demo

前言

使用OpenCL获得GPU的能力

从手机上导出动态库并在工程中使用

使用dlopen

cl_stub

扩展

实现在GPU进行计算

clsl定义

定龚俊义测试数据

创建环境

数据写入

执行

关于工作组和工作项的设置

读取运算结approve果

扩展-使用OpenGL实现计算

扩展-获取设备GPUGitHub基本信息

最后

作者信息

推广

在Android上实现Metal的计算Demo

前言

使用OpenCL获得GPU的能力

从手机上导出动态库并在工程中使用

使用dlopen

cl_stub

扩展

实现在GPU进行计算

clsl定义

定龚俊义测试数据

创建环境

数据写入

执行

关于工作组和工作项的设置

读取运算结approve果

扩展-使用OpenGL实现计算

扩展-获取设备GPUGitHub基本信息

最后

相关文章

多分辨率适配的思考

实战编程·使用SwiftUI从0到1完成一款iOS笔记App（五）

Systrace和RheaTrace使用

iOS 性能优化 — APP启动时间解析

作者信息

推广