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)