Skip to content

Instantly share code, notes, and snippets.

@ykst
Created January 28, 2015 14:05
  • Star 4 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
Star You must be signed in to star a gist
Save ykst/c1b38bbea93df28da3ed to your computer and use it in GitHub Desktop.
iOSでもOpenCLを使えるか調べてみた

iOSでもOpenCLを使えるか調べてみた

いつまで経ってもPrivate Frameworksに引き蘢っているiOSのOpenCL APIは果たして本当に使えるのか? というわけでアプリから引き摺り出してみます。Jailbreakはしません。

とりあえずdlopen(3)してみる

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)で動的にシンボル解決をすれば実機でも使える筈です。

OpenCLバージョンを表示してみる

まずはスモークテストとしてバージョンを表示してみます。

#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_NAMECL_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の登場で大分楽にはなりましたし、そんなわけで粛々とフラグメントシェーダをハックする作業に戻るとしましょう。

Work In Progress

まだひょっとしたら抜け道があるかもしれませんが、 検証用に使ったXcodeプロジェクトを晒しておくので、我こそはという人は自己責任でチャレンジしてみてください。

@ykst
Copy link
Author

ykst commented Sep 27, 2015

Metalがkernelをサポートしてくれているので、OpenCLはこのまま闇に消し去られるだろう

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment