月別アーカイブ: 2012年1月

RenderScript Tegra 3 の Quad core で Multi Thread 実行

RenderScript で Android 3.x でも Multi Thread 実行する方法が
わかりました。
以下その測定結果

         最小実行時間 速度比  OS       CPU
----------------------------------------------------------
Tegra 3        622ms   x2.6   A3.2.1   Cortex-A9 1.3GHz x4
OMAP4460      1334ms   x1.2   A4.0.1   Cortex-A9 1.2GHz x2
Tegra 2       1608ms   x1.0   A3.1     Cortex-A9 1.0GHz x2

 * 実行時間の数値が小さいほうが速い。
 * 何度か実行して最速時間を抽出。

(Eee Pad TF201, Galaxy Nexus SC-04D, Optimus Pad L-06C)

やっと RenderScript で Quad core による効果を確認できるように
なりました。さすがに Tegra 3 の CPU は速いです。

前の記事で書いたとおり Android 3.x の HelloCompute サンプルは
シングルスレッドで実行されます。

マルチスレッド実行できない条件は下記の通り

 RenderScript で自分自身のカーネルを呼び出す rsForEach() を記述する。

つまり下記のようなコードを用いて、rsForEach() が同じファイルの
root() 関数である (A) を呼び出すとだめです。

// RenderScript

void root( const uchar4 *v_in, uchar4 *v_out ) // -- (A)
{
  ~
}

rs_script  script;
uchar4*	   in_vec;
uchar4*	   out_vec;

void filter()
{
    rsForEach( script, rsGetAllocation( in_vec ), rsGetAllocation( out_vec ), 0 );
}

これを回避する方法は 2つあります。

●(1) Android 4.0 (API 14) 以降のみ

前回の記事に書いたとおり Java から直接 forEach_root() を使って
呼び出します。
Android 4.0 (API Level 14) 以降の HelloCompute はこれです。
同じ RenderScript 内の rsForEach() 呼び出しを通らないので実行が
スレッド並列化されます。

●(2) Android 3.0 (API 11) 以降対応

別の RenderScript から rsForEach() を使って呼び出します。

(2) の方法を使ったマルチスレッド実行できる RenderScript の例
(Android 3.2 Sample の HelloCompute を修正)

// RenderScript: mono.rs

#pragma version(1)
#pragma rs java_package_name(com.example.android.rs.hellocompute)

// 実際に実行したい関数
void root(const uchar4 *v_in, uchar4 *v_out, const void *usrData, uint32_t x, uint32_t y)
{
    float4 f4 = rsUnpackColor8888(*v_in);
    float3 mono = dot(f4.rgb, gMonoMult);
    *v_out = rsPackColorTo8888(mono);
}

ダミーの RenderScript ファイルを作ります。call.rs

// RenderScript: call.rs

#pragma version(1)
#pragma rs java_package_name(com.example.android.rs.hellocompute)

void root(const uchar4 *v_in, uchar4 *v_out)
{ // ダミー関数。呼ばれない。
}

rs_script script;
uchar4*	  in_vec;
uchar4*	  out_vec;

void callscript()
{
    rsForEach( script, rsGetAllocation( in_vec ), rsGetAllocation( out_vec ), 0 );
}

call.rs 経由で mono.rs を呼び出します。

// Java: HelloCompute.java

    mScript = new ScriptC_mono(mRS, getResources(), R.raw.mono);

    // 追加
    mCall = new ScriptC_call(mRS, getResources(), R.raw.call);

    mCall.bind_in_vec( mInAllocation );
    mCall.bind_out_vec( mOutAllocation );
    mCall.set_script( mScript );

    // 削除
    //mScript.set_gIn(mInAllocation);
    //mScript.set_gOut(mOutAllocation);
    //mScript.set_gScript(mScript);
    //mScript.invoke_filter();

    // call.rs 経由で呼び出す
    mCall.invoke_callscript();

    mOutAllocation.copyTo(mBitmapOut);

●マルチスレッド実行の検出

mono.rs にわざと競合が発生するような修正を加えます。

// RenderScript: mono.rs

static uint32_t sum= 0;
static uint32_t sum2= 0;

void root(const uchar4 *v_in, uchar4 *v_out)
{
    uint32_t  a= sum;

    float4 f4 = rsUnpackColor8888(*v_in);
    float3 mono = dot(f4.rgb, gMonoMult);
    *v_out = rsPackColorTo8888(mono);

    sum= a+1;
    rsAtomicInc( &sum2 ); // Android 4.0 API 14 以降のみ
}

void dump_sum()
{
    rsDebug( "sum=", sum );
    rsDebug( "sum2=", sum2 );
}

実行後に Java から dump_sum() を呼び出します。

// Java: HelloCompute.java

    mCall.invoke_callscript();
    mOutAllocation.copyTo(mBitmapOut);

    mScript.invoke_dump_sum();  // 結果表示

期待通り sum が毎回不定な値となり、並列実行されている証拠が
取れました。

Android 4.0 以降では rsAtomic 命令が使えます。
上の例の場合 sum2 は常に一定となり、呼ばれた回数 == 画素数
を示すようになりました。

同一ファイル内の rsForEach() では、カーネルコード実行後に
static な変数にアクセスできます。
上の例のように、競合する集約されたコードが記述されている可能性があり
互換性の問題があったのかもしれません。

forEach のエントリは常に root() で、それぞれ個別のファイルに記述する
必要があります。C言語でも独立性が保たれます。
並列実行しやすいようこのような仕様になっているものと考えられます。

関連エントリ
Android 4.0 RenderScript と MultiCore CPU
Android 3.x RenderScript (7) RenderScript Compute の速度
Android 4.0 RenderScript Compute の速度 その2

Android 4.0 RenderScript と MultiCore CPU

以前 Tegra2 や Galaxy Nexus (OMAP4460) など Dual core CPU
搭載機種で RenderScript の動作テストを行いました。

 ・Android 3.x RenderScript (7) RenderScript Compute の速度
 ・Android 4.0 RenderScript Compute の速度 その2

予想に反してあまり速くなく、Single thread で書いた NDK コードと
ほとんど変わらない結果でした。
また Atomic 命令なしに競合するコードを記述しても不定な結果になりません。
そのため RenderScript は Single thread 実行しているのではないか
との結論でした。

その後 Eee Pad TF201 (Android 3.2) の Quad core CPU Tegra 3 で
走らせても同じで、Single Thread 相当の速度しか出ません。
使い方に問題があるかもしれないので、いろいろと試したところ
rsForEach() と forEach_root() で速度が違うことがわかりました。

API Level 14 (Android OS 4.0) 以降は RenderScript の interface が
変更されており、java から直接 cl の root を呼び出すことができます。

下記は API Level による HelloCompute のサンプルの違いです。

// RenderScript : API 11-13 (Android 3.x)

#pragma version(1)
#pragma rs java_package_name(com.example.android.rs.hellocompute)

rs_allocation gIn;
rs_allocation gOut;
rs_script gScript;

const static float3 gMonoMult = {0.299f, 0.587f, 0.114f};

void root(const uchar4 *v_in, uchar4 *v_out, const void *usrData, uint32_t x, uint32_t y) {
    float4 f4 = rsUnpackColor8888(*v_in);

    float3 mono = dot(f4.rgb, gMonoMult);
    *v_out = rsPackColorTo8888(mono);
}

void filter() {
    rsForEach(gScript, gIn, gOut, 0);
}

Android 4.0 以降は filter() がありません。

// RenderScript : API 14-15 (Android 4.0.x)

#pragma version(1)
#pragma rs java_package_name(com.example.android.rs.hellocompute)

const static float3 gMonoMult = {0.299f, 0.587f, 0.114f};

void root(const uchar4 *v_in, uchar4 *v_out) {
    float4 f4 = rsUnpackColor8888(*v_in);

    float3 mono = dot(f4.rgb, gMonoMult);
    *v_out = rsPackColorTo8888(mono);
}

Java から invoke_filter() 経由で script 内で rsForEach() を呼び出す
必要がなく、直接 forEach_root() を使って root() を呼べるからです。

// Java : API 11-13 (Android 3.x)

        mScript = new ScriptC_mono(mRS, getResources(), R.raw.mono);

        mScript.set_gIn(mInAllocation);
        mScript.set_gOut(mOutAllocation);
        mScript.set_gScript(mScript);
        mScript.invoke_filter();
        mOutAllocation.copyTo(mBitmapOut);

Allocation の受け渡しもなく簡単になっています。

// Java : API 14-15 (Android 4.0.x)

        mScript = new ScriptC_mono(mRS, getResources(), R.raw.mono);

        mScript.forEach_root(mInAllocation, mOutAllocation);
        mOutAllocation.copyTo(mBitmapOut);

Galaxy Nexus SC-04D (Android 4.0.1 API Level 14) でこの両者の
速度を比べると、API Level 14-15 の forEach_root() を用いた方が
およそ 2倍の速度で実行されます。
dual core による並列実行と考えられます。

更に調べると、rs 内の外部参照可能な関数に rsForEach() を使った
呼び出しを記述するだけで速度が遅くなることがわかりました。
つまり一番上の「// RenderScript : API 11-13 (Android 3.x)」を
forEach_root() で実行しても 1/2 の速度となります。
互換性のためなのか、条件によって thread 実行できない script と
みなされる何らかの理由があるのだと考えられます。

Android 3.x のデバイスでは forEach_root() は使えませんでした。
Tegra 3 など Quad core CPU の本当の能力が発揮されるのは
Android 4.0 へ アップデートしてからかもしれません。

続く RenderScript Tegra 3 の Quad core で Multi Thread 実行

関連エントリ
Android 3.x RenderScript (7) RenderScript Compute の速度
Android 4.0 RenderScript Compute の速度 その2

ASUS Eee Pad TF201 Tegra3 の GPU 速度

Transformer Prime, EeePad TF201 手に入れました。
NVIDIA の新しい Mobile プロセッサ Tegra 3 を搭載しています。
CPU は Cortex-A9 の quad core で GPU は ULP GeForce 12 core 版
となります。

CPU は 4 core 認識しています。Tegra2 と違い NEON あります。

Processor	: ARMv7 Processor rev 9 (v7l)
processor	: 0
BogoMIPS	: 1992.29

processor	: 1
BogoMIPS	: 1992.29

processor	: 2
BogoMIPS	: 1992.29

processor	: 3
BogoMIPS	: 1992.29

Features	: swp half thumb fastmult vfp edsp neon vfpv3 
CPU implementer	: 0x41
CPU architecture: 7
CPU variant	: 0x2
CPU part	: 0xc09
CPU revision	: 9

GPU の方は Tegra2 と比べて特に機能的な違いがありませんでした。
depth は 16bit のままで depth_texture もありません。

Mobile GPU の機能比較

GL_VERSION: OpenGL ES 2.0 12.01014
GL_RENDERER: NVIDIA Tegra 3
GL_VENDOR: NVIDIA Corporation
GL_SHADING_LANGUAGE_VERSION: OpenGL ES GLSL 1.00

Extension は下記のページに追加しました。

OpenGL ES Extension

● GPU 速度

いつもの pixel 負荷の高い描画テストを走らせてみました。
Tegra2 と比べて 3倍弱といった数値になっています。

(2) light 3 (ambient + directional x2 + point )
-----------------------------------------------------
Tegra 3 / ULP GeForce(12)  15.70fps   EeePad TF201
Tegra 2 / ULP GeForce(8)    5.74fps   ICONIA TAB A500


(3) light 1 (ambient + directional )
-----------------------------------------------------
Tegra 3 / ULP GeForce(12)  30.10fps   EeePad TF201
Tegra 2 / ULP GeForce(8)   12.10fps   ICONIA TAB A500

shader core は Pixel が 2倍らしいので、Mali400 と同じ表現を
使えば MP2 相当です。2倍以上の性能差は動作クロックの違いだと
考えられます。

他の GPU との比較は下記のページにまとめました。

Mobile GPU bench mark

pix/sec の欄を見ると、速度的にはほぼ 3世代目グループである
Mali-400MP4/Adreno220/SGX543MP2 に並んでいることがわかります。

NVIDIA といえば GPU のイメージが強いですが Tegra はむしろ逆です。
優先しているのは CPU の方で、半世代先行した設計で Quad 化を
実現しています。
それに対して GPU は決して突出した性能ではなく、やっとライバルに
並んだところです。

Group3 | Adreno 220, PVR SGX543MP2, Mali-400MP4, ULP GeForce(12,Tegra3)
Group2 | Adreno 205, PVR SGX530/535/540, ULP GeForce(8,Tegra250)
Group1 | Adreno 200, AMD Z430

関連エントリ
OpenGL ES 2.0 Mobile GPU の比較、重いテクスチャ
頂点性能の比較 その2 (OpenGL ES 2.0 Mobile GPU)
OpenGL ES 2.0 Mobile GPU の頂点性能を比較する
A5 PowerVR SGX543MP2 は iOS 5 だと速い
さらに OpenGL ES 2.0 Mobile GPU の速度比較
OpenGL ES 2.0 Mobile GPU の速度比較 (dual core世代) 更新

Android 4.0 対応、RenderScript のサンプル

昨年の 11月頃に掲載していた RenderScript のソースコードが
Galaxy Nexus Android 4.0 で動作しなかったので修正しました。

flatlib_ap02a.zip
flatlib_ap02b.zip
flatlib_ap02c.zip
flatlib_ap02d.zip

Constant (Uniform) の初期値を Java から書き込んでいる部分が
原因でした。RenderScript 上で書きこまなければならなかった
ようで、そのように修正しています。

例えば ap02d では RenderScript 側で初期化用の関数を作成し
Java から invoke で呼び出しています。

// RenderScript
~
void setup()
{
    vconst->CameraPosition.x= 0.0f;
    vconst->CameraPosition.y= 0.0f;
    vconst->CameraPosition.z= 3.0f;
    vconst->CameraPosition.w= 0.0f;
    ~
}

最後に初期化コードを呼び出しておく。

// Java
    ~
    mScript.set_rstate( rstate );
    mScript.bind_vconst( geometry_const );
    mScript.bind_mconst( material_const );
    mScript.bind_light( light_const );

    mScript.invoke_setup(); // 初期化関数の呼び出し

    mRS.bindRootScript( mScript );

関連エントリ
Android 3.x RenderScript で 3D モデルのライティング
Android 3.x RenderScript (5) 任意の 3D モデルを描画する
Android 3.x RenderScript (4) script で頂点を書き換える
Android 3.x RenderScript (3) 独自シェーダーの割り当てとメッシュの描画(2D)

Android 4.0 RenderScript Compute の速度 その2

Galaxy Nexus を使う機会があったので RenderScript を走らせてみました。
Android 4.0 搭載かつ TI OMAP4460 は Cortex-A9 dual に NEON が
搭載されています。

[Tegra2] Optimus Pad L-06C   | Android 3.1   API Level 11
 Tegra250 1.0GHz (Cortex-A9 dual) VFPv3-D16 (NEON無し)

[4460]   Galaxy Nexus SC-04D | Android 4.0.1 API Level 14
 OMAP4460 1.2GHz (Cortex-A9 dual) VFPv3-D32 + NEON

                                      [Tegra2] [4460]
                                        VFP     VFP  NEON
--------------------------------------------------------------
(1) RenderScript rsMatrixMultiply       161     111   116 msec
(2) RenderScript 乗算展開               248     176   185
(3) RenderScript vec命令展開            162     111   118
(4) Java 乗算展開                       564     417   424
(5) Java android.opengl.Matrix 呼出    2536    8864  8861
(6) NDK C++ 関数呼び出し                170     104    95
(7) NDK C++ 乗算展開                    166      65    98
(8) NDK C++ 乗算展開v2                  128      58    66
(9) NDK C++ NEON intrinsic               --      --    61
(10)NDK inline asm fmacs                118      58    59
(11)NDK inline asm NEON                  --      --    52
(12)NDK inline asm fmacs + pthread x2    76      56    50
(13)NDK inline asm NEON + pthread x2     --      --    48

(実行時間、数値は msec。値が小さいほうが高速。5回の平均値)

NDK の実行バイナリは VFP (-mfpu=vfp) と NEON (-mfpu=neon) の
2種類あります。Tegra2 は VFP 版だけ走らせています。
NDK 部分だけの違いで、RenderScript や Java の結果には影響しません。

(3) は rsMatrixMultiply() 相当を float4 の積和に置き換えたものです。
(1) とほぼ同等の速度が得られています。

// RenderScript (3)
rs_matrix4x4	TransformMatrix;

void root( const float4* vin, float4* vout )
{
    float4  tmp;
    tmp=  *(float4*)&TransformMatrix.m[ 0] * vin->x;
    tmp+= *(float4*)&TransformMatrix.m[ 4] * vin->y;
    tmp+= *(float4*)&TransformMatrix.m[ 8] * vin->z;
    tmp+= *(float4*)&TransformMatrix.m[12] * vin->w;
    *vout= tmp;
}

(8) は NDK でも (3) 同様に最適化されやすい書き方に変更しています。
やはりこちらの方が効率が良いようです。
ただし生成コードを追った限りでは特にベクタ化されておらず
neon 命令が含まれていませんでした。

// NDK: C++ (8)
static void loopndk21( const Vect4* vin, Vect4* vout, int length )
{
    Matrix  tmat= TransformMatrix.Transpose();

    for( int i= 0 ; i< length ; i++ ){
        Vect4   tmp;

        tmp.x=  tmat._11 * vin->x;
        tmp.y=  tmat._21 * vin->x;
        tmp.z=  tmat._31 * vin->x;
        tmp.w=  tmat._41 * vin->x;

        tmp.x+= tmat._12 * vin->y;
        tmp.y+= tmat._22 * vin->y;
        tmp.z+= tmat._32 * vin->y;
        tmp.w+= tmat._42 * vin->y;

        tmp.x+= tmat._13 * vin->z;
        tmp.y+= tmat._23 * vin->z;
        tmp.z+= tmat._33 * vin->z;
        tmp.w+= tmat._43 * vin->z;

        tmp.x+= tmat._14 * vin->w;
        tmp.y+= tmat._24 * vin->w;
        tmp.z+= tmat._34 * vin->w;
        tmp.w+= tmat._44 * vin->w;

        *vout++= tmp;
        vin++;
    }
}

(9) は (8) を元にコンパイラの neon 命令で記述したものです。
実行速度にはあまり大きな差が出ませんでした。
演算が単純なのでそれ以外の部分がボトルネックになっている
可能性があります。

// NDK: C++ (9)
static void loopndk5( const Vect4* vin, Vect4* vout, int length )
{
    Matrix  tmat= TransformMatrix.Transpose();
    float32x4_t*    fout= (float32x4_t*)vout;
    float32x4_t     m0= *(float32x4_t*)&tmat._11;
    float32x4_t     m1= *(float32x4_t*)&tmat._12;
    float32x4_t     m2= *(float32x4_t*)&tmat._13;
    float32x4_t     m3= *(float32x4_t*)&tmat._14;
    for( int i= 0 ; i< length ; i++ ){
        float32x4_t tmp;
        tmp= vmulq_n_f32(      m0, vin->x );
        tmp= vmlaq_n_f32( tmp, m1, vin->y );
        tmp= vmlaq_n_f32( tmp, m2, vin->z );
        tmp= vmlaq_n_f32( tmp, m3, vin->w );
        *fout++= tmp;
        vin++;
    }
}

(9) と同じ物をインラインアセンブラで書いたものです。
neon は非常に少ない命令で記述することが可能です。

// NDK: inline asm (11)/(13)
    __asm__ __volatile( "\
        vld4.32 {d0,d2,d4,d6},[%0]!\n\
        vld4.32 {d1,d3,d5,d7},[%0]\n\
    1: \n\
        vldmia %1!, {d8-d9} \n\
        vmul.f32 q5,q0,d8[0] \n\
        vmla.f32 q5,q1,d8[1] \n\
        vmla.f32 q5,q2,d9[0] \n\
        vmla.f32 q5,q3,d9[1] \n\
        vstmia %2!, {d10-d11} \n\
        subs  %3,%3,#1 \n\
        bne 1b \n\
    "
    : "=&r"( mat ), "=&r"( vin ), "=&r"( vout ), "=&r"( length )
    : "0"( mat ), "1"( vin ), "2"( vout ), "3"( length )
    :
        "d0", "d1", "d2", "d3", "d4", "d5", "d6", "d7",
        "d8", "d9", "d10", "d11", "cc"
    );

OMAP4460 は CPU クロックで Tegra250 より 1.2倍速い計算になりますが、
結果はそれ以上の差が開いています。
演算能力よりもバス速度など他の原因が考えられます。

C/asm の結果でも neon は vfp よりわずかに速いだけで大きな差が
生じませんでした。
RenderScript がどの程度ハードウエアの能力を引き出しているのか
判断するにはまだデータが足りないようです。

関連エントリ
Android 3.x RenderScript (7) RenderScript Compute の速度
Android 3.x RenderScript (6) Graphics まとめ
Android 3.x RenderScript (5) 任意の 3D モデルを描画する
Android 3.x RenderScript (4) script で頂点を書き換える
Android 3.x RenderScript (3) 独自シェーダーの割り当てとメッシュの描画(2D)
Android 3.x RenderScript (2) 描画と Allocation
Android 3.x RenderScript (1)