いつまで経ってもPrivate Frameworksに引き蘢っているiOSのOpenCL APIは果たして本当に使えるのか? というわけでアプリから引き摺り出してみます。Jailbreakはしません。
Private Frameworksは基本的にSDK配下のものと同じパスでランタイムにdlopen(3)
する事が出来るので、実機で開いてみます。
#import "AppDelegate.h"
#import <dlfcn.h>
- (BOOL)application:(UIApplication *)application didFinishLaunchingWithOptions:(NSDictionary *)launchOptions {
void *lib = dlopen("/System/Library/PrivateFrameworks/OpenCL.framework/OpenCL", RTLD_LAZY);
printf("%p\n", lib);
return YES;
}
iOS6,7,8それぞれで見つかりました。ここまでは簡単です。
APIはどの程度サポートされているのか?ということでnm(1)
でライブラリシンボルを探ってみます。
OSXのnm(1)
には --defined-only
が無いのでcomm(1)
で外部ライブラリシンボルとの差分を取ってみます。
# 以下のパスはXcodeのインストール先に依存
$ cd /Applications/Xcode.app/Contents/Developer/Platforms/iPhoneOS.platform/Developer/SDKs/iPhoneOS8.1.sdk/System/Library/PrivateFrameworks/OpenCL.framework
$ comm -23 <(nm -gj OpenCL | sort) <(nm -uj OpenCL | sort) | uniq
ずらずらっとOpenCL APIのシンボルが出てきます。
OSXのOpenCL.frameworkにあるcl*.h
辺りはちゃんと定義されていたので、
そこの定義を借用してdlsym(3)
で動的にシンボル解決をすれば実機でも使える筈です。
まずはスモークテストとしてバージョンを表示してみます。
#include <OpenCL/cl.h> // 適当にパスを通してAPI宣言を関数ポインタのexternに変更しておく
#include <stdio.h>
#include <dlfcn.h>
CL_API_ENTRY cl_int CL_API_CALL
(*clGetPlatformIDs)(cl_uint /* num_entries */,
cl_platform_id * /* platforms */,
cl_uint * /* num_platforms */) CL_API_SUFFIX__VERSION_1_0;
CL_API_ENTRY cl_int CL_API_CALL
(*clGetPlatformInfo)(cl_platform_id /* platform */,
cl_platform_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
static __attribute__((constructor)) void
cl_bridge_load_symbols(void)
{
void *lib = dlopen("/System/Library/PrivateFrameworks/OpenCL.framework/OpenCL", RTLD_LAZY);
clGetPlatformIDs = dlsym(lib, "clGetPlatformIDs");
clGetPlatformInfo = dlsym(lib, "clGetPlatformInfo");
cl_uint num = 0;
clGetPlatformIDs(0, 0, &num);
cl_platform_id platforms[num];
clGetPlatformIDs(num, platforms, &num);
char version[1024] = {};
clGetPlatformInfo (platforms[0], CL_PLATFORM_VERSION,
1024, version, 0);
version[1023] = '\0';
printf("Version: %s\n", version);
}
iPhone5S+iOS8.1だと以下のように出ます。
Version: OpenCL 1.2 (Sep 12 2014 13:36:29)
なんだかいけそうです。 後はこの調子でAPIを宣言していけば普通にOpenCLを使う事が出来る、 そんなふうに考えていた時期が僕にもありました。
結論から言うと__ほぼまともに使う事は出来ませんでした__。 試した条件とどこまで使えたかを下表にまとめてみます。
デバイス / iOS | APIバージョン | Device *1 | clGetPlatformInfo | clGetDeviceInfo | clBuildProgram | clEnqueueNDRangeKernel |
---|---|---|---|---|---|---|
3GS / 6.1.6 | 1.1 | CPU | ○ | ○ | ○ | △ *2 |
iPad3 / 7.1.2 | 1.2 | CPU | ○ | ○ | △ *3 | × |
5S / 8.1 | 1.2 | GPU | ○ | × | × | × |
*1) CL_DEVICE_NAME
とCL_DEVICE_TYPE_(CPU|GPU)
から判断
*2) local work sizeに1以上を設定出来ない
*3) コンパイラは反応するが、CL_BUILD_PROGRAM_FAILURE
(エラーログ無し)
iPad3/iOS7ではデバイスアクセスまでは出来るのですが、コンパイルを通してもCL_SUCCESS
してくれないため、そこより先に進む事は出来ませんでした。
5S/iOS8に至ってはclGetDeviceInfo
の段階でEXC_BAD_ACCESS
が出る系が有り、ほぼ何も出来ませんでした。IOKit経由でGPUに問い合わせをする経路がまとめて遮断されているような気がします。
というわけで、唯一実際にカーネルを走らせる事が出来たのが3GS/iOS6のみ、
それもCPU実装で、local work sizeの合計数に1より大きい値を設定するとCL_INVALID_WORK_GROUP_SIZE
が返ってくるというなんとも寒い結果と相成りました。
基本的にiOSはバージョンが上がる毎にPrivate APIのガードが固くなっていくため、まだiOS6やiOS7.0では何とか使えるパターンがあったのかもしれません。試していませんが、rootを取れればまた話は違うでしょう。 普通に使えたらエンタープライズ用途で何か応用出来るかなと思っていたのですが、もう僕の心は折れました。
やはりiOSでGPGPUをしたければ今まで通りglslをexploitするしかなさそうです。 ES3の登場で大分楽にはなりましたし、そんなわけで粛々とフラグメントシェーダをハックする作業に戻るとしましょう。
まだひょっとしたら抜け道があるかもしれませんが、 検証用に使ったXcodeプロジェクトを晒しておくので、我こそはという人は自己責任でチャレンジしてみてください。
Metalがkernelをサポートしてくれているので、OpenCLはこのまま闇に消し去られるだろう