この記事は自分用の備忘録です。解説用ではないです
Android NDKを使って c/c++のコードは書けるけど、OpenCLはどこから手を付けていいかわからない人（つまり私！）の参考になれば幸いです。

きっかけ

カメラから取得した画像を、ゴニョゴニョっと画像処理して表示する、というアプリを開発しています。

約1年前にc++とアセンブリ言語で画像を反転させてみるという記事を書き、その後、自分の必要な処理に関するアセンブリ言語くらいは書けるようになり、実行速度が大幅にアップしました。

アセンブリ言語でSIMDに対応したコードを書いても、ある程度行くと処理速度に限界点があり、さらなる進化をしていくためには「GPUプログラミングしかない」、となったので、以前から興味のあったOpenCLに手を出すことにしました。

ちなみに「ハードウェアをアップグレードする」という選択肢が選択できないのが辛いところです。
逆に言うと、ハードウェアは一択なので、とりあえずそれで動けば今回はOKです。

OpenCL

OpenCLとは、現時点の理解度でするとマルチコアCPU・GPU上でコードを動かすための共通言語ということみたいです。

各GPUメーカーがそれぞれ特色を持ったGPUを出しています。
それぞれSDKを提供しているのですが、それぞれ命令セットや使い方が異なるので、それを共通した仕様のC言語で使えるようにしようと策定したのがOpenCLだそう。

GPUだけでなくCPU上でも動くとのこと。

OpenCLに入門するに当たり、何冊かまとめて購入した中で、一番初心者に向いていると思って現在読んでいる本はこれです。

改訂新版 OpenCL入門 1.2対応マルチコアCPU・GPUのための並列プログラミング

なんかレビューを見ると、「誰かの参考になるとは思えない」とか「対応年数は過ぎたと思います」とか、酷評されてますねｗ
とりあえず私の役には立ってます。

ホストとデバイス

OpenCLには2つの立場があるそうです。それが「ホスト」と「デバイス」。
大雑把な説明で言えば下記の通り。

ホスト・・・・メインコードが動いているCPU
デバイス・・・コードを動かすためのGPU

先程「GPUでコードを動かすための共通言語」と書きましたが、もう少し正確には「GPUでコードを動かしてもらうための共通フレームワーク」のようです。
CPUからGPUに「こうゆうコードを動かしてよ(=カーネルコード)」とお願いするので、CPU側がホスト、GPU側をデバイスと呼びます。
ただしホストCPUをデバイスとして使うこともできるそうです。

「はじめてのOpelCLプログラム」にチャンレンジ

目指すところはGPU上で画像処理なのですが、とりあえず最初のコードとしては「Hello world」と表示する、とか、そのくらいのレベルでないと辛いです。
しかもネット上にある資料の殆どは、Windows, Linux, Macで動かすことを想定して書いているので、Android用に懇切丁寧に解説してくれている記事は、殆どなかったです。
とりあえず参考書のHello worldを書いてみました。

はじめてのOpelCLプログラム (hello.cl)

__kernel _void hello(__global char* string)
{
  string[0] = 'H';
  string[1] = 'e';
  string[2] = 'l';
  string[3] = 'l';
  string[4] = 'o';
  string[5] = ',';
  string[6] = ' ';
  string[7] = 'W';
  string[8] = 'o';
  string[9] = 'r';
  string[10] = 'l';
  string[11] = 'd';
  string[12] = '!';
  string[13] = '\0';
}

はじめてのOpelCLプログラム (helloworld.c)

#include <stdio.h>
#include <stdlib.h>
#include <CL/cl.h>

#define MEM_SIZE (128)
#define MAX_SOURCE_SIZE (0x100000)

int main()
{
  cl_device_id device_id = NULL;
  cl_context context = NULL;
  cl_command_queue command_queue = NULL;
  cl_mem memobj = NULL;
  cl_program program = NULL;
  cl_kernel kernel = NULL;
  cl_platform_id platform_id = NULL;
  cl_uint ret_num_devices;
  cl_uint ret_num_platforms;
  cl_int ret;

  char string[MEMSIZE];
  FILE *fp;
  char fileName[] = "./hello.cl";
  char *source_str;
  size_t source_size;

  /* カーネルを含むソースコードをロード */
  fp = fopen(fileName, "r");
  if (!fp) {
    fprintf(stderr, "Failed to load kernel.\n");
    exit(1);
  }
  source_str = (char*)malloc(MAX_SOURCE_SIZE);
  source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
  fclose(fp);

  /* プラットフォーム・デバイスの情報取得 */
  ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
  ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);

  /* OpenCLコンテキストの作成 */
  context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);

  /* コマンドキューの作成 */
  command_queue = clCreateCommandQueue(context, device_id, 0, &ret);

  /* メモリバッファの作成 */
  memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(char), NULL, &ret);

  /* 読み込んだソースからカーネルプログラムをビルド */
  program = clCreateProgramWithSource(context, 1, (const char**)&source_str, (const size_t *)&source_size, &ret);

  /* カーネルプログラムをビルド */
  ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

  /* OpenCLカーネルの作成 */
  kernel = clCreateKernel(program, "hello", &ret);

  /* OpenCLカーネル引数の設定 */
  ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);

  /* OpenCLカーネルを実行 */
  ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);

  /* メモリバッファから結果を取得 */
  ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(char), string, 0, NULL, NULL);

  /* 結果の表示 */
  puts(string);

  /* 終了処理 */
  ret = clFlush(command_queue);
  ret = clFinish(command_queue);
  ret = clReleaseKernel(kernel);
  ret = clReleaseProgram(program);
  ret = clReleaseMemObject(memobj);
  ret = clReleaseCommandQueue(command_queue);
  ret = clReleaseContext(context);

  free(source_str);

  return 0;

Android用のはどこ・・・？

Android NDKの中には、CL/cl.hがありません。いろいろ探し回って、https://github.com/krrishnarraj/libopencl-stub が使えました。

Android用にコードを修正

上記コードをAndroid用に修正します。ファイルツリーはこんな感じ。

main.h

#include <jni.h>
#include "helloworld.h"

extern "C" {

  void helloworld(JNIEnv* env, jobject self);

  // http://kznote.blogspot.com/2013/01/androidnative.html
  //
  // B - byte
  // C - char
  // D - double
  // F - float
  // I - int
  // J - long
  // S - short
  // V - void
  // Z - boolean
  // l - jobject
  // L(java class name); - Java class name. i.e : Ljava/lang/String;
  // [(class name) - array of class name.  i.e.: byte str[]  -> [B
  static JNINativeMethod jniMethodTable[] = {
      {"native_helloworld", "()V", (void *) helloworld}
  };

  JNIEXPORT jint JNI_OnLoad(JavaVM* vm, void* reserved);
}

main.cpp

#include "main.h"

// Register c++ functions to VM
JNIEXPORT jint JNI_OnLoad(JavaVM *vm, void *reserved) {

  // Register export functions to JAVA
  JNIEnv *env;
  if (vm->GetEnv((void **) &env, JNI_VERSION_1_6) != JNI_OK) {
    return JNI_ERR;
  } else {
    jclass clazz = env->FindClass("test/wf9a5m75/opencl/MainActivity");
    if (clazz) {
      int method_table_size = sizeof(jniMethodTable) / sizeof(jniMethodTable[0]);
      jint ret = env->RegisterNatives(clazz, jniMethodTable, method_table_size);
      env->DeleteLocalRef(clazz);
      return ret == 0 ? JNI_VERSION_1_6 : JNI_ERR;
    } else {
      return JNI_ERR;
    }
  }

  return JNI_VERSION_1_6;
}

helloworld.h

#include <jni.h>
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <android/log.h>

#define MY_APP_LOG_TAG "opencl_test"
#define LOGD(...) __android_log_print(ANDROID_LOG_DEBUG, MY_APP_LOG_TAG, __VA_ARGS__)
#define LOGE(...) __android_log_print(ANDROID_LOG_ERROR, MY_APP_LOG_TAG, __VA_ARGS__)

helloworld.c

#include "helloworld.h"
#include "hello_cl.h"

#define MEM_SIZE (128)
#define MAX_SOURCE_SIZE (0x100000)

void helloworld(JNIEnv* env, jobject self)
{
  cl_device_id device_id = NULL;
  cl_context context = NULL;
  cl_command_queue command_queue = NULL;
  cl_mem memobj = NULL;
  cl_program program = NULL;
  cl_kernel kernel = NULL;
  cl_platform_id platform_id = NULL;
  cl_uint ret_num_devices;
  cl_uint ret_num_platforms;
  cl_int ret;

  char txt[MEM_SIZE];


  size_t source_size;

  /* カーネルを含むソースコードをロード */
  source_size = sizeof(HELLO_CL);

  /*
   * プラットフォーム・デバイスの情報取得
   *   OpenCLが動作するプラットフォーム(=ハードウェア)を認識して
   *   該当するプラットフォームを第2引数(platform_id)に返す。
   *   以降はこのplatform_idを参照することで、プラットフォームが利用できる
   *
   * args[0] = ホスト側のアプリケーションが望むプラットフォームの数(通常は1)
   * args[1] = プラットフォームIDが返される
   * args[2] = 実際に利用できるOpenCL対応のプラットフォームの数が返される
   */
  ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
  if (ret == CL_SUCCESS) {
    LOGE("--->clGetPlatformIDs = success");
    LOGE("       ret_num_platforms = %d", ret_num_platforms);
    LOGE("       platform_id = 0x%08x", (int)platform_id);
  } else {
    print_error("clGetPlatformIDs", ret);
  }

  /*
   * デバイスの特定
   *   プラットフォームで使うGPUといったデバイスを特定する。
   *
   * args[0] = プラットフォームID
   * args[1] = デバイスの種類を指定
   *   CL_DEVICE_TYPE_DEFAULT : 標準デバイス
   *   CL_DEVICE_TYPE_GPU : 明示的にGPU
   *   CL_DEVICE_TYPE_CPU : ホスト側のCPUをデバイスとして使用
   * args[2] = 利用したいデバイスの数
   * args[3] = デバイスIDが返される
   * args[4] = args[1]で指定したデバイスの数か返される。利用できるデバイスがないときは0
   */
  ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
  if (ret == CL_SUCCESS) {
    LOGE("--->clGetDeviceIDs = success");
  } else {
    print_error("clGetDeviceIDs", ret);
  }

  /*
   * OpenCLコンテキストの作成
   *   OpenCLを動作させる実行環境となるOpenCLコンテキストを作成する。
   *   以降作成する各種OpenCLのオブジェクトは、このコンテキストに所属することになり、
   *   同一コンテキスト内のオブジェクトを通して各デバイスを制御できる。
   *   このコンテキストは演算デバイスが1つ以上利用できる1つの仮想コンピュータと捉えると良いらしい。
   *
   * args[0] = ?
   * args[1] = 利用するデバイスの数
   * args[2] = 利用するデバイスに対応するデバイスハンドルのリスト
   * args[3] = ?
   * args[4] = ?
   * args[5] = 戻り値
   *
   * ret = OpenCLコンテキスト
   */
  context =  clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
  if (ret == CL_SUCCESS) {
    LOGE("--->clCreateContext = success");
  } else {
    print_error("clGetDeviceIDs", ret);
  }

  /*
   * コマンドキューの作成
   *   OpenCLではコマンドキューを通して、ホストからデバイスに対する働きかけ
   *   (カーネル実行コマンドや、ホストーデバイス間のメモリコピーコマンド）を
   *   実行する。1つのデバイスに対して必ず1つ以上のコマンドキューを作成する。
   *
   * args[0] = コンテキスト
   * args[1] = デバイスID
   * args[2] = ?
   * args[3] = 戻り値
   *
   * ret = コマンドキュー
   */
  command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
  if (ret == CL_SUCCESS) {
    LOGE("--->clCreateCommandQueue = success");
  } else {
    print_error("clCreateCommandQueue", ret);
  }

  /*
   * メモリオブジェクトの作成
   *   デバイス側のメモリにホスト側からアクセスするためのメモリオブジェクトの作成。
   *   デバイスはデバイス内のメモリにしかアクセスできないので、ホスト側が
   *   デバイス内のメモリに必要なデータを設定する。
   *
   * args[0] = コンテキスト
   * args[1] = メモリ属性
   * args[2] = 確保するメモリ容量
   * args[3] = ?
   * args[4] = 戻り値
   *
   * ret = メモリオブジェクト
   */
  memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(char), NULL, &ret);
  if (ret == CL_SUCCESS) {
    LOGE("--->clCreateBuffer = success");
  } else {
    print_error("clCreateBuffer", ret);
  }

  /*
   * 読み込んだソースコードからカーネルプログラムを作成
   *   OpenCLではカーネルプログラムをまずプログラムオブジェクトとして認識する。
   *   コンパイル済バイナリから作成する場合は、clCreateProgramWithBinary()を利用する。
   *
   * args[0] = コンテキスト
   * args[1] = ?
   * args[2] = カーネルプログラムのソースコード
   * args[3] = 同、バイト数
   */
  program = clCreateProgramWithSource(context, 1, &HELLO_CL, NULL, &ret);
  if (ret == CL_SUCCESS) {
    LOGE("--->clCreateProgramWithSource = success");
  } else {
    print_error("clCreateProgramWithSource", ret);
  }

  /*
   * カーネルプログラムをビルド
   *   プログラムオブジェクトをOpenCL Cコンパイラ・リンカを使用してビルド。
   *
   * args[0] = プログラムオブジェクト
   * args[1] = args[2]で指定するデバイスの数
   * args[2] = デバイスリストへのポインタ
   * args[3] = コンパイラオプションの文字列
   * args[4] = ?
   * args[5] = ?
   */
  ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
  if (ret == CL_SUCCESS) {
    LOGE("--->clBuildProgram = success");
  } else if (ret == CL_DEVICE_NOT_FOUND) {
    print_error("clBuildProgram", ret);
  }

  /*
   * OpenCLカーネルの作成
   *   カーネルオブジェクトの作成。
   *   1つのカーネルオブジェクトは1つのカーネル関数に対応するので
   *   作成時に関数名を作成する。
   *
   *   複数のカーネル関数を1つにプログラムオブジェクトとして記述してもOK。
   *   ただし1つのカーネルオブジェクトは1つのカーネル関数に1対1なので、clCreateKernel()を複数回呼び出すことになる。
   *
   * args[0] = プログラムオブジェクト
   * args[1] = 関数名
   * args[2] = 戻り値
   *
   * ret = カーネルオブジェクト
   */
  kernel = clCreateKernel(program, "hello", &ret);
  if (ret == CL_SUCCESS) {
    LOGE("--->clCreateKernel = success");
  } else {
    print_error("clCreateKernel", ret);
  }

  /*
   * OpenCLカーネル引数の設定
   *   hello()関数のargs[0] stringがどこにあるのかを指定する。
   *   __global付きカーネル引数に対しては、
   *   ホスト側から確保したデバイスメモリを示すメモリオブジェクトを割り当てる。
   *
   * args[0] = カーネルオブジェクト
   * args[1] = カーネル引数の位置
   * args[2] = args[3]のサイズ
   * args[3] = 与える値(この例ではメモリオブジェクト)
   *
   * ----------------------
   * ホスト側のデータをカーネル引数として直接値渡しする場合の例
   *   int a = 10;
   *   clSetKernelArg(kernel, 0, sizeof(int), (void *)&a);
   */
  ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);
  if (ret == CL_SUCCESS) {
    LOGE("--->clSetKernelArg = success");
  } else {
    print_error("clSetKernelArg", ret);
  }

  /*
   * OpenCLカーネルを実行(コマンドキューに投入)
   *   デバイスにカーネル関数hello()を実行するようにリクエストする。
   *   この関数が終了した時点ではコマンドキューに投入されただけで
   *   実行が完了したことは保証されない。
   *   完了を待つためには、args[4]でイベントオブジェクトを取得する必要があるが
   *   HelloWorldの時点では紹介されていないので、省略。
   *
   * args[0] = コマンドキュー
   * args[1] = 投入するカーネル
   * args[2] = ?
   * args[3] = ?
   * args[4] = ?
   */
  ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);
  if (ret == CL_SUCCESS) {
    LOGE("--->clEnqueueTask = success");
  } else {
    print_error("clEnqueueTask", ret);
  }

  /*
   * メモリオブジェクトを通じて結果を取得
   *   デバイス側のメモリからホスト側のメモリにデータをコピーする。
   *   (逆の場合はclEnqueueWriteBuffer())
   *
   *   Enqueueと名前にある通り、キューを通じて実行される
   *
   * args[0] = コマンドキュー
   * args[1] = デバイス側のメモリ
   * args[2] = データ完了まで待つかどうか
   *   CL_TRUE : 待つ(同期コピー)
   *   CL_FALSE : 待たない(非同期コピー)
   * args[3] = args[4]のコピーするサイズ
   * args[4] = コピー先(ホスト側のメモリのポインタ)
   * args[5] = オフセット?
   * args[6] = ?
   * args[7] = ?
   */
  ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(char), txt, 0, NULL, NULL);
  if (ret == CL_SUCCESS) {
    LOGE("--->clEnqueueReadBuffer = success");
  } else {
    print_error("clEnqueueReadBuffer", ret);
  }


  /*
   * 結果の表示
   */
  LOGD("answer = %s", txt);

  /*
   * 終了処理
   */
  ret = clFlush(command_queue);
  ret = clFinish(command_queue);
  ret = clReleaseKernel(kernel);
  ret = clReleaseProgram(program);
  ret = clReleaseMemObject(memobj);
  ret = clReleaseCommandQueue(command_queue);
  ret = clReleaseContext(context);
}


void print_error(char *name, cl_int ret) {
  if (ret == CL_DEVICE_NOT_FOUND) {
    LOGE("--->%s = CL_DEVICE_NOT_FOUND", name);
  } else if (ret == CL_DEVICE_NOT_AVAILABLE) {
    LOGE("--->%s = CL_DEVICE_NOT_AVAILABLE", name);
  } else if (ret == CL_COMPILER_NOT_AVAILABLE) {
    LOGE("--->%s = CL_COMPILER_NOT_AVAILABLE", name);
  } else if (ret == CL_MEM_OBJECT_ALLOCATION_FAILURE) {
    LOGE("--->%s = CL_MEM_OBJECT_ALLOCATION_FAILURE", name);
  } else if (ret == CL_OUT_OF_RESOURCES) {
    LOGE("--->%s = CL_OUT_OF_RESOURCES", name);
  } else if (ret == CL_OUT_OF_HOST_MEMORY) {
    LOGE("--->%s = CL_OUT_OF_HOST_MEMORY", name);
  } else if (ret == CL_PROFILING_INFO_NOT_AVAILABLE) {
    LOGE("--->%s = CL_PROFILING_INFO_NOT_AVAILABLE", name);
  } else if (ret == CL_MEM_COPY_OVERLAP) {
    LOGE("--->%s = CL_MEM_COPY_OVERLAP", name);
  } else if (ret == CL_IMAGE_FORMAT_MISMATCH) {
    LOGE("--->%s = CL_IMAGE_FORMAT_MISMATCH", name);
  } else if (ret == CL_IMAGE_FORMAT_NOT_SUPPORTED) {
    LOGE("--->%s = CL_IMAGE_FORMAT_NOT_SUPPORTED", name);
  } else if (ret == CL_BUILD_PROGRAM_FAILURE) {
    LOGE("--->%s = CL_BUILD_PROGRAM_FAILURE", name);
  } else if (ret == CL_MAP_FAILURE) {
    LOGE("--->%s = CL_MAP_FAILURE", name);
  } else if (ret == CL_MISALIGNED_SUB_BUFFER_OFFSET) {
    LOGE("--->%s = CL_MISALIGNED_SUB_BUFFER_OFFSET", name);
  } else if (ret == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST) {
    LOGE("--->%s = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST", name);
  } else if (ret == CL_INVALID_VALUE) {
    LOGE("--->%s = CL_INVALID_VALUE", name);
  } else if (ret == CL_INVALID_DEVICE_TYPE) {
    LOGE("--->%s = CL_INVALID_DEVICE_TYPE", name);
  } else if (ret == CL_INVALID_PLATFORM) {
    LOGE("--->%s = CL_INVALID_PLATFORM", name);
  } else if (ret == CL_INVALID_DEVICE) {
    LOGE("--->%s = CL_INVALID_DEVICE", name);
  } else if (ret == CL_INVALID_CONTEXT) {
    LOGE("--->%s = CL_INVALID_CONTEXT", name);
  } else if (ret == CL_INVALID_QUEUE_PROPERTIES) {
    LOGE("--->%s = CL_INVALID_QUEUE_PROPERTIES", name);
  } else if (ret == CL_INVALID_COMMAND_QUEUE) {
    LOGE("--->%s = CL_INVALID_COMMAND_QUEUE", name);
  } else if (ret == CL_INVALID_HOST_PTR) {
    LOGE("--->%s = CL_INVALID_HOST_PTR", name);
  } else if (ret == CL_INVALID_MEM_OBJECT) {
    LOGE("--->%s = CL_INVALID_MEM_OBJECT", name);
  } else if (ret == CL_INVALID_IMAGE_FORMAT_DESCRIPTOR) {
    LOGE("--->%s = CL_INVALID_IMAGE_FORMAT_DESCRIPTOR", name);
  } else if (ret == CL_INVALID_IMAGE_SIZE) {
    LOGE("--->%s = CL_INVALID_IMAGE_SIZE", name);
  } else if (ret == CL_INVALID_SAMPLER) {
    LOGE("--->%s = CL_INVALID_SAMPLER", name);
  } else if (ret == CL_INVALID_BINARY) {
    LOGE("--->%s = CL_INVALID_BINARY", name);
  } else if (ret == CL_INVALID_BUILD_OPTIONS) {
    LOGE("--->%s = CL_INVALID_BUILD_OPTIONS", name);
  } else if (ret == CL_INVALID_PROGRAM) {
    LOGE("--->%s = CL_INVALID_PROGRAM", name);
  } else if (ret == CL_INVALID_PROGRAM_EXECUTABLE) {
    LOGE("--->%s = CL_INVALID_PROGRAM_EXECUTABLE", name);
  } else if (ret == CL_INVALID_KERNEL_NAME) {
    LOGE("--->%s = CL_INVALID_KERNEL_NAME", name);
  } else if (ret == CL_INVALID_KERNEL_DEFINITION) {
    LOGE("--->%s = CL_INVALID_KERNEL_DEFINITION", name);
  } else if (ret == CL_INVALID_KERNEL) {
    LOGE("--->%s = CL_INVALID_KERNEL", name);
  } else if (ret == CL_INVALID_ARG_INDEX) {
    LOGE("--->%s = CL_INVALID_ARG_INDEX", name);
  } else if (ret == CL_INVALID_ARG_VALUE) {
    LOGE("--->%s = CL_INVALID_ARG_VALUE", name);
  } else if (ret == CL_INVALID_ARG_SIZE) {
    LOGE("--->%s = CL_INVALID_ARG_SIZE", name);
  } else if (ret == CL_INVALID_KERNEL_ARGS) {
    LOGE("--->%s = CL_INVALID_KERNEL_ARGS", name);
  } else if (ret == CL_INVALID_WORK_DIMENSION) {
    LOGE("--->%s = CL_INVALID_WORK_DIMENSION", name);
  } else if (ret == CL_INVALID_WORK_GROUP_SIZE) {
    LOGE("--->%s = CL_INVALID_WORK_GROUP_SIZE", name);
  } else if (ret == CL_INVALID_WORK_ITEM_SIZE) {
    LOGE("--->%s = CL_INVALID_WORK_ITEM_SIZE", name);
  } else if (ret == CL_INVALID_GLOBAL_OFFSET) {
    LOGE("--->%s = CL_INVALID_GLOBAL_OFFSET", name);
  } else if (ret == CL_INVALID_EVENT_WAIT_LIST) {
    LOGE("--->%s = CL_INVALID_EVENT_WAIT_LIST", name);
  } else if (ret == CL_INVALID_EVENT) {
    LOGE("--->%s = CL_INVALID_EVENT", name);
  } else if (ret == CL_INVALID_OPERATION) {
    LOGE("--->%s = CL_INVALID_OPERATION", name);
  } else if (ret == CL_INVALID_GL_OBJECT) {
    LOGE("--->%s = CL_INVALID_GL_OBJECT", name);
  } else if (ret == CL_INVALID_BUFFER_SIZE) {
    LOGE("--->%s = CL_INVALID_BUFFER_SIZE", name);
  } else if (ret == CL_INVALID_MIP_LEVEL) {
    LOGE("--->%s = CL_INVALID_MIP_LEVEL", name);
  } else if (ret == CL_INVALID_GLOBAL_WORK_SIZE) {
    LOGE("--->%s = CL_INVALID_GLOBAL_WORK_SIZE", name);
  } else if (ret == CL_INVALID_PROPERTY) {
    LOGE("--->%s = CL_INVALID_PROPERTY", name);
  } else {
    LOGE("--->%s is failed", name);
  }
}

hello_cl.h

#ifndef _HELLO_CL_
const char* HELLO_CL =
    "__kernel void hello(__global char* string)\n"
    "{\n"
    "  string[0] = 'H';\n"
    "  string[1] = 'e';\n"
    "  string[2] = 'l';\n"
    "  string[3] = 'l';\n"
    "  string[4] = 'o';\n"
    "  string[5] = ',';\n"
    "  string[6] = ' ';\n"
    "  string[7] = 'W';\n"
    "  string[8] = 'o';\n"
    "  string[9] = 'r';\n"
    "  string[10] = 'l';\n"
    "  string[11] = 'd';\n"
    "  string[12] = '!';\n"
    "  string[13] = '\\0';\n"
    "}";

#endif

Android.mk

LOCAL_PATH := $(call my-dir)

include $(CLEAR_VARS)

LOCAL_MODULE := libhelloworld
LOCAL_SRC_FILES :=  main.cpp helloworld.c
LOCAL_STATIC_LIBRARIES += OpenCL
LOCAL_ARM_MODE := arm
LOCAL_C_INCLUDES += $(LOCAL_PATH)/../libopencl-stub/include
LOCAL_LDLIBS := -llog

include $(BUILD_SHARED_LIBRARY)

#libOpenCL.soはどこ？

今回実験に使っているのは、Firefly-RK3288という評価ボード。
本番環境のSoCがRK3288なので、この評価ボードで試してます。

RK3288はRockchip社が製造するシリーズのCPUで、OpenCL 1.1が使えることになってます。

http://opensource.rock-chips.com/wiki_RK3288

さて、先程のコードを実行すると、私の環境では動きませんでした。

「なんでやねん」と思いながら調べて、ようやく見つけました。

The problem is some android devices don't have opencl lib file libopencl.so in file system, or the file has a different name(for example, libgles_Mali.so).
https://stackoverflow.com/a/29427228

まじっすか！

ということで、findコマンドで探すと、あっ。あった。

ということで、このファイルのパスをlibopencl-stubに追加して、再度実行！

ｷﾀ━━━━(ﾟ∀ﾟ)━━━━!!

Github repo

ここに置いておきます。
https://github.com/wf9a5m75/opencl_test