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

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)