在Android上實現Metal的計算Demo

語言: CN / TW / HK

theme: cyanosis

前言

Metal的官方文件中有一個利用GPU進行計算的例子:Performing Calculations on a GPU。本文主要探討的是如何在Android上實現類似的效果

使用OpenCL獲得GPU的能力

之所以選用OpenCL,是因為它的api比較符合通用計算的場景。雖然它不僅適合GPU,還適合其他單元。但本文主要會利用OpenCL來進行GPU運算

GPU 優化技術-OpenCL 介紹

由於NDK並沒有提供關於OpenCL的動態庫,所以在Android中使用OpenCL需要弄一些騷操作。

從手機上匯出動態庫並在工程中使用

以筆者使用的測試機Oppo R15為例,OpenCL環境可在裝置的/system/vendor/lib/libOpenCL.so獲取。此時可通過adb pull將so包匯出,放入工程中使用。

bash adb pull /system/vendor/lib/libOpenCL.so 還有一個比較重要的是CL的標頭檔案,這個可以在官方的github上下載:

需要注意的是: 1. 由於OpenCL的版本問題,上述標頭檔案定義的某些邏輯,對應到您在裝置上匯出的so包不一定試用。需要按OpenCL真實的版本進行程式碼編寫。 2. 由於so包是在單一裝置匯出的,不排除在其他裝置有適配問題,且一般高通平臺Adreno gpu命名為libOpenCL.so,但其他平臺則不然。

此部分可參考: 高通平臺下安卓opencl小例子

使用dlopen

筆者希望能夠儘可能多的在不同裝置上執行,所以想盡量使用在裝置上的環境。

參考csarron/cl_stub: Ease OpenCL library loading可知,我們可以通過dlopen來獲取OpenCL環境以及特定方法的控制代碼。簡單分析一下該庫的做法。

cl_stub

首先還是需要將OpenCL的標頭檔案匯入到專案中: - c的標頭檔案:KhronosGroup/OpenCL-Headers - c++的封裝:KhronosGroup/OpenCL-CLHPP

定義一個可能的OpenCL動態庫路徑陣列: c 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環境時,迴圈遍歷該陣列,尋找一個可以讀取的路徑。如果可以讀取則被認為是可用於載入的屬於該裝置的OpenCL動態庫路徑。 ```c 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之後,會有一個嘗試將path的路徑替換成apk下路徑的操作,目的是為了使用自己apk的so包,但該操作存在一些問題會導致dlopen返回NULL`。建議將此段程式碼註釋*。

最後通過dlopen獲取該庫的例項,用於後面方法呼叫的使用。 c so_handle = dlopen(path, RTLD_LAZY);

在該庫中也有實現與OpenCL的api名字相同的函式,我們以獲取平臺的api-clGetPlatformIDs為例: ```c 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;
}

} `` 這裡通過上述事例出來的例項so_handle,**利用dlsym獲取到clGetPlatformIDs`的控制代碼,繼而呼叫它。至此就完成了一次呼叫OpenCL的api對接**。

ps:f_clGetPlatformIDs為定義的一個與clGetPlatformIDs引數結構一致的函式別名 c 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的做法,alibaba/MNN使用OpenCL能力的程式碼中也有體現

MNN/OpenCLWrapper.cpp有類似邏輯:

```c++ static const std::vector 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(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++封裝,您也可以直接使用其c程式碼的api效果是一致的。

clsl定義

在Metal的計算例子中,我們可以看到它定義的核心函式: c++ 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中,我們也能定義一個clslc++ 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(result)的每項。gid和Metal中的index定義都為該執行緒在網格中的位置。這裡可將其表示為第幾次呼叫該函式進行運算。

定義測試資料

```c++ 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); } `` 這裡我們按照Metal的Demo,給資料的長度設定為1 << 24,併為datadata2`填入隨機數。

建立環境

```c++ std::vector 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 devices = context.getInfo(); 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陣列可根據不同平臺設定特殊的常量。這一塊由於沒有深入研究就不多細說了,詳細可參考[MNN/OpenCLRuntime.cpp](https://github.com/alibaba/MNN/blob/master/source/backend/opencl/core/runtime/OpenCLRuntime.cpp)。 - 根據上下文獲取裝置context.getInfo(),一般陣列只有一項,取第0位即可。 - 根據上下文和裝置建立一個CommandQueue,用於後續的函式執行。 - 利用上述clsl,建立一個Program,並對clsl進行編譯。 -Kernelclsl中函式的對映,譬如上述定義的add方法,則會被例項成Kernel`。

資料寫入

```c++ 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); `` 上述邏輯將之前造的資料(data、data2、result)建立對應gpu能使用的Buffer。並將其設定到Kernel中**作為函式的輸入**。對應的是clsl`中的三個輸入。

CL_MEM_COPY_HOST_PTR等flag的使用會關係到該記憶體在哪裡建立的問題。這裡的定義會在device中開闢一塊記憶體(即所謂的視訊記憶體),並將原有的記憶體(即所謂的主存)賦值給他(data、data2、result)。這裡會發生一次記憶體拷貝,屬於耗時操作

詳細定義可參考:從零開始學習OpenCL開發(三)深入API

執行

```c++ 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(); `` 將Kernel新增到CommandQueue中,由此可見gpu的執行對於cpu是非同步操作。最後可通過queue.finish();`同步等待執行完成。

關於工作組和工作項的設定

工作組和工作項在Metal中也可定義為執行緒組和執行緒queue.enqueueNDRangeKernel需要設定全部工作項數(global_work_size)和每個工作組的工作項數(local_work_size)。

  • 一個工作組可容納多少個工作項,這個是有最大值的,可通過Kernel獲取。即上述的workgroupSize
  • 全部工作項數,因為這裡我們的運算可以理解為迴圈相加1 << 24,所以可以設定為1 << 24

所以最後網格中有1 << 24個工作項,(1 << 24)/ workgroupSize個工作組。因為這裡的運算只是一維的,所以只需考慮一維即可。

此部分可參考:

OpenCl 對維度,工作項,工作組概念的理解

OpenCL與Metal API下如何合理地安排執行緒組大小

Metal框架詳細解析(十五) —— 計算處理之關於執行緒和執行緒組(三)

擴充套件閱讀:

OpenCL優化:工作組大小效能優化

讀取運算結果

c++ queue.enqueueReadBuffer(cBuffer, CL_TRUE, 0, bufferSize, result); 從deivce中將運算結果拷貝回result,作為結果的輸出讀取。

至此,在Android上利用GPU進行運算的效果就基本實現了。Demo會貼在文末,以供參考。

擴充套件-使用OpenGL實現計算

OpenGL也是可以實現簡單的數值運算的,只是由於它本身不是設計來支援簡單的運算的。所以不能很好的適配。在OpenGL ES 3.0後可以利用Transform Feedback在頂點著色器運算好的結果對映並取出。具體流程如下:

我們可以定義一個用於計算的頂點著色器和一個空的片元著色器: ```c++ 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 FeedbackoutValue,讀取出計算結果。這裡可以使用EGL搭建OpenGL的環境

  • 詳細可參考文章:

OpenGL - Transform Feedback

OpenGL ES 3.0使用Transform Feedback進行通用計算並讀取結果

Android OpenGL ES 系列連載:(07)Transform Feedback ps:這裡面的Demo有關於EGL的封裝

擴充套件-獲取裝置GPU基本資訊

利用OpenCL可獲取裝置GPU的基本資訊,這裡列舉幾個最基本的,譬如deviceNamedeviceVersion。 ```c++ std::vector 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(); LOGCATD("cl_platform_name:%s", platformName.c_str());

std::vector 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(); LOGCATD("cl_device_name:%s", deviceName.c_str()); const std::string deviceVersion = gpuDevices[0].getInfo(); LOGCATD("cl_device_version:%s", deviceVersion.c_str()); const std::string deviceVendor = gpuDevices[0].getInfo(); 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上搜索得到,安裝後如果裝置支援OpenCL,可以在上面看到詳細的裝置資訊以及不同平臺對於它的擴充套件介面,以及庫在系統的路徑

最後

本文主要介紹如何利用OpenCL在Android上實現Metal的計算邏輯,以及相關的環境適配與擴充套件。最後貼出上述的Demo,共同學習!ps:筆者使用的是ndk版本16.1.4479499,cmake版本3.6.4111459,可按需修改。