D3D Shader/OpenGL」カテゴリーアーカイブ

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)

Android 3.x RenderScript で 3D モデルのライティング

RenderScript で光源やカメラ位置を使った描画方法を取り上げて欲しいとの
メールをいただいたので試してみました。

renderscript_ap02d.jpg

OpenGL ES 2.0 なのでこれらのパラメータの反映はシェーダー任せです。
演算は RenderScript よりも GLSL の範疇になるかと思いますが、
複数の ConstantBuffer の作製や PixelShader への渡し方などを参考にして
いただければと思います。

●必要なデータ構造の定義と初期化

RenderScript で、必要な情報の構造体を宣言します。
内容は Java, RenderScript, Shader でそれぞれ共有します。

ジオメトリ関連

// RenderScript
typedef struct GeometryConstant {
    rs_matrix4x4    Projection;		// Projection Matrix
    rs_matrix4x4    World;		// World Matrix
    rs_matrix4x4    PView;		// RS → Shader 用
    float4          CameraPosition;	// (x y z -) カメラ位置
    float4          CameraAim;		// (x y z -) 注視点
} GeometryConstant_T;

描画オブジェクトのマテリアル

// RenderScript
typedef struct MaterialConstant {
    float4      AmbientColor;	// (r g b -)
    float4      DiffuseColor;	// (r g b -)
    float4      SpecularColor;	// (r g b p)
} MaterialConstant_T;

光源。Intensity は Color に事前乗算しておく。

// RenderScript
typedef struct LightConstant {
    float4      AmbientLightColor;		// (r g b -)
    float4      DirectionalLightColor;		// (r g b -)
    float4      PointLightColor;		// (r g b -)
    float4      DirectionalLightDirection;	// (x y z -)
    float4      PointLightPosition;		// (x y z -)
} LightConstant_T;

Java 上でインスタンスを作って初期化します。
カメラは座標と注視点を登録しています。

// Java
// Geometry Uniform
ScriptField_GeometryConstant  geometry_const= new ScriptField_GeometryConstant( mRS, 1, Allocation.USAGE_SCRIPT|Allocation.USAGE_GRAPHICS_CONSTANTS );

// Projection Matrix
Matrix4f    projection= new Matrix4f();
projection.loadPerspective( 30.0f, mWidth/mHeight, 0.1f, 100.0f );
geometry_const.set_Projection( 0, projection, true );

// Camera Position
geometry_const.set_CameraPosition( 0, new Float4( 0.0f, 0.0f, 3.0f, 0.0f ), true );
geometry_const.set_CameraAim( 0, new Float4( 0.0f, 0.0f, 0.0f, 0.0f ), true );

// World Matrix
Matrix4f    world= new Matrix4f();
world.loadIdentity();
geometry_const.set_World( 0, world, true );

テクスチャを使わないのでマテリアルカラーを設定。

// Java
// Material Uniform
ScriptField_MaterialConstant  material_const= new ScriptField_MaterialConstant( mRS, 1, Allocation.USAGE_SCRIPT|Allocation.USAGE_GRAPHICS_CONSTANTS );

material_const.set_AmbientColor( 0, new Float4( 0.3f, 0.3f, 0.3f, 1.0f ), true );
material_const.set_DiffuseColor( 0, new Float4( 1.0f, 0.9f, 0.9f, 1.0f ), true );
material_const.set_SpecularColor( 0, new Float4( 1.0f, 1.0f, 0.4f, 30.0f ), true );

光源を設定

// Light Uniform
ScriptField_LightConstant  light_const= new ScriptField_LightConstant( mRS, 1, Allocation.USAGE_SCRIPT|Allocation.USAGE_GRAPHICS_CONSTANTS );

light_const.set_AmbientLightColor( 0, new Float4( 0.3f, 0.3f, 0.3f, 1.0f ), true );
light_const.set_DirectionalLightColor( 0, new Float4( 1.0f, 1.0f, 1.0f, 1.0f ), true );
light_const.set_PointLightColor( 0, new Float4( 0.0f, 1.0f, 1.0f, 1.0f ), true );
light_const.set_DirectionalLightDirection( 0, new Float4( 0.57f, -0.57f, -0.57f, 0.0f ), true );
light_const.set_PointLightPosition( 0, new Float4( 0.0f, 0.2f, 0.0f, 0.0f ), true );

作成した Constant Buffer をシェーダーで読めるようにします。
Builder で必要な数だけ addConstant() して Uniform を予約します。
インスタンスは ProgramVertex の bindConstants() で渡します。

必要な constnat buffer だけ渡しています。
VertexShader は GeometryConstant と LightConstant。

// Java
// VertexShader
ProgramVertex.Builder   vsh_builder= new ProgramVertex.Builder( mRS );
vsh_builder.setShader( mRes, R.raw.mesh3d_vsh );
vsh_builder.addConstant( geometry_const.getType() );
vsh_builder.addConstant( light_const.getType() );
vsh_builder.addInput( model.getElement() );

ProgramVertex   vsh= vsh_builder.create();
vsh.bindConstants( geometry_const.getAllocation(), 0 );
vsh.bindConstants( light_const.getAllocation(), 1 );

PixelShader (FragmentShader) も同様です。
MaterialConstant/LightConstant の 2つです。

// Java
// PixelShader
ProgramFragment.Builder psh_builder= new ProgramFragment.Builder( mRS );
psh_builder.setShader( mRes, R.raw.mesh3d_psh );
psh_builder.addConstant( material_const.getType() );
psh_builder.addConstant( light_const.getType() );

ProgramFragment psh= psh_builder.create();
psh.bindConstants( material_const.getAllocation(), 0 );
psh.bindConstants( light_const.getAllocation(), 1 );

Shader と同じように、RenderScript にもアクセスするデータを渡します。
GeometryConstant/LightConstant の 2つ。

// Java
mScript.bind_vconst( geometry_const );
mScript.bind_light( light_const );

●RenderScript

ViewMatrix (WorldSpace To CameraSpace) の作成と ProjectionMatrix
との前計算を行なっています。

CameraPosition/CameraAim, ProjectionMatrix を元に、
PView (ProjectionMatrix * ViewMatrix) を求めます。

// RenderScript
rs_matrix4x4    viewMatrix;
LookAt( &viewMatrix, vconst->CameraPosition.xyz, vconst->CameraAim.xyz );
rsMatrixLoadMultiply( &vconst->PView, &vconst->Projection, &viewMatrix );

RenderScript は C言語ながら、上記のようにシェーダーのような vector 記法
が使えます。例えば float4 で宣言した変数も .xyz をつけることで float3
の部分アクセスができます。

float4 vec;
float3 pos= vec.xyz;

swizzle もできました。使いやすいです。

float4 d= vec.wzyx;
float4 e= vec.yyyy;
float4 f;
f.zw= vec.xy;

RenderScritp に LookAt 関数があるかと思ったらなかったので作ります。

// RenderScript
static void LookAt( rs_matrix4x4* view, float3 pos, float3 aim )
{
    rsMatrixLoadIdentity( view );
    float3  up= { 0.0f, 1.0f, 0.0f };
    float3  pz= normalize( aim - pos );
    float3  px= normalize( cross( pz, up ) );
    float3  py= normalize( cross( px, pz ) );
    view->m[0]= px.x;
    view->m[4]= px.y;
    view->m[8]= px.z;
    view->m[1]= py.x;
    view->m[5]= py.y;
    view->m[9]= py.z;
    view->m[2]= -pz.x;
    view->m[6]= -pz.y;
    view->m[10]= -pz.z;
    rsMatrixTranslate( view, -pos.x, -pos.y, -pos.z );
}

これ以外に、サンプルの script では効果がわかりやすいように
カメラ位置と点光源のアニメーションを行なっています。

●Shader

シェーダーでは与えられたパラメータを元に座標と色の計算を行なっています。
ScreenSpace だけでなく WorldSpace での座標と法線の算出。

// GLSL
varying vec3    onormal;
varying vec3    oeye;
varying vec3    opos;

void main()
{
    vec4    wpos= UNI_World * vec4( ATTRIB_position.xyz, 1.0 );
    gl_Position= UNI_PView * wpos;
    onormal.xyz= mat3(UNI_World) * ATTRIB_normal.xyz;
    oeye.xyz= UNI_CameraPosition.xyz - wpos.xyz;
    opos.xyz= wpos.xyz;
}

MaterialConstant と LightConstant を用いてライティング。

// GLSL
precision mediump   float;

varying vec3 onormal;
varying vec3 oeye;
varying vec3 opos;

void main()
{
    vec3    normal= normalize( onormal.xyz );

    // AmbientLight
    lowp vec3    diffuse= UNI_AmbientLightColor.xyz * UNI_AmbientColor.xyz;

    // DirectinalLight
    float   dd= clamp( dot( normal.xyz, -UNI_DirectionalLightDirection.xyz ), 0.0, 1.0 );
    diffuse.xyz+= UNI_DirectionalLightColor.xyz * UNI_DiffuseColor.xyz * dd;

    float   ds= pow( clamp( dot( normal.xyz, normalize( oeye.xyz - UNI_DirectionalLightDirection.xyz ) ), 0.0, 1.0 ), UNI_SpecularColor.w );
    lowp vec3    specular= UNI_SpecularColor.xyz * ds;

    // PointLight
    vec3    lightvec= UNI_PointLightPosition.xyz - opos.xyz;
    vec3    lightdir= normalize( lightvec.xyz );
    float   dl= dot( lightvec.xyz, lightvec.xyz ) * 10.0;
    float   pd= clamp( dot( normal.xyz, lightdir.xyz ), 0.0, 1.0 ) / dl;
    diffuse.xyz+= UNI_PointLightColor.xyz * UNI_DiffuseColor.xyz * pd;

    gl_FragColor.xyz= diffuse.xyz + specular.xyz;
    gl_FragColor.w= 1.0;
}

●ソースコード

モデルデータの読み込みは以前のエントリと同じ物を使っています。

flatlib_ap02d.zip ソースコード

●最適化

無駄が多いので最適化の余地があります。

ConstantBuffer は更新するデータ量が少なくなるように分割できます。
今回は Java, RenderScript, Shader でバッファを共有していますが不要な
パラメータがあります。RenderScript で書き換えるものと Shader だけが
参照するものをわけると転送量を減らせます。

それでも ConstantBuffer (Uniform) は頂点書き換えと異なり PushBuffer
を経由するため lock が発生しないため効率は良いです。

PixelShader (FragmentShader) での演算は負荷が重いので、出来る限り
シェーダーの外に出した方が良いです。
MaterialColor * LightColor はプリミティブ単位で求まるので前計算可能です。
頂点単位で十分なものは VertexShader に分散させることができます。

関連エントリ
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)

Android 3.x RenderScript (7) RenderScript Compute の速度

RenderScript の用途は大きく分けて二通りあります。
Compute と Graphics です。

Compute (rs_cl) は rsForEach() 命令を使い、CUDA/OpenCL のような
並列処理を意識した呼び出し方となっています。
おそらくマルチコア CPU による並列化や GPU が将来的にサポートされる
のではないかと考えられます。

Graphics (rs_graphics) は少々特殊で、描画ループを Native 実行する
C言語で記述できることが特徴です。今までにないレイヤーですが、
NDK の部分的な置き換えを狙っているようにみえます。
またシームレスに rs_cl を呼び出せるのも大きなメリットです。

RenderScript は他にも vector や matrix 演算がサポートされており、
NEON 等の SIMD 命令に展開されるのではないかと考えられます。
ただし現状だと Android 3.x (Honeycomb) デバイスの大半が
NEON 命令が使えない NVIDIA Tegra 250 (Tegra2) なので、
その効果を確認することができませんでした。

●速度テスト

RenderScript は本当に速いのか、実際に速度を調べてみました。
Matrix4x4 * float4 を 100万頂点分演算します。

Android 3.1 API Level 11
LG Optimus Pad L-06C
Tegra250 1GHz (Cortex-A9 dual) NEON 無し VFPv3-D16

(1) RenderScript rsMatrixMultiply      160msec
(2) RenderScript 乗算展開              248msec
(3) Java 乗算展開                      571msec
(4) Java android.opengl.Matrix 呼出   2537msec
(5) NDK C++ 関数呼び出し               167msec
(6) NDK C++ 乗算展開                   168msec
(7) NDK inline asm fmacs               117msec
(8) NDK inline asm fmacs + pthread      72msec

(実行時間。値が小さいほうが高速)

データの転送時間は含まれていません。演算とループのみです。
ただし (8) に関してはスレッドの起動時間、終了同期(join) の時間が
含まれています。

●RenderScript

(1) は確かに速いですが、普通に NDK と C++ だけで書いたコード
(5)/(6) とほとんど差がありませんでした。
NEON があればもっと大きく差がついたのかもしれません。

同時に dual core CPU であっても、RenderScript は特に並列実行
していないことがわかります。

(2) のように自分で乗算展開した場合はかなり遅くなっています。
NDK の Cコンパイラよりも通常演算の最適化がまだ弱いようです。

// RenderScript: (1)
rs_matrix4x4    TransformMatrix;
void root( const float4* vin, float4* vout )
{
    *vout= rsMatrixMultiply( &TransformMatrix, *vin );
}

// RenderScript: (2)
rs_matrix4x4    TransformMatrix;
void root( const float4* vin, float4* vout )
{
    vout->x= TransformMatrix.m[ 0] * vin->x
            +TransformMatrix.m[ 4] * vin->y
            +TransformMatrix.m[ 8] * vin->z
            +TransformMatrix.m[12] * vin->w;

    vout->y= TransformMatrix.m[ 1] * vin->x
            +TransformMatrix.m[ 5] * vin->y
            +TransformMatrix.m[ 9] * vin->z
            +TransformMatrix.m[13] * vin->w;

    vout->z= TransformMatrix.m[ 2] * vin->x
            +TransformMatrix.m[ 6] * vin->y
            +TransformMatrix.m[10] * vin->z
            +TransformMatrix.m[14] * vin->w;

    vout->w= TransformMatrix.m[ 3] * vin->x
            +TransformMatrix.m[ 7] * vin->y
            +TransformMatrix.m[11] * vin->z
            +TransformMatrix.m[15] * vin->w;
}

●Java

(3) が予想よりも高速でした。
演算自体の負荷が高い場合は比較的言語の差が出にくいのですが、
JIT が効果的に働いているのではないかと思われます。
結果を Java で受け取るならば、データ転送時間が不要な分もっと
差が縮まる可能性があります。

// java : (3)
float[] fv= transformMatrix.getArray();
float[] destbuf= new float[MAX_TRANSFORM*4];

long    stime= System.currentTimeMillis();
for( int i= 0 ; i< MAX_TRANSFORM ; i++ ){
    int base= i*4;
    destbuf[base+0]=
          fv[ 0] * srcbuf[base+0]
        + fv[ 4] * srcbuf[base+1]
        + fv[ 8] * srcbuf[base+2]
        + fv[12] * srcbuf[base+3];

    destbuf[base+1]=
          fv[ 1] * srcbuf[base+0]
        + fv[ 5] * srcbuf[base+1]
        + fv[ 9] * srcbuf[base+2]
        + fv[13] * srcbuf[base+3];

    destbuf[base+2]=
          fv[ 2] * srcbuf[base+0]
        + fv[ 6] * srcbuf[base+1]
        + fv[10] * srcbuf[base+2]
        + fv[14] * srcbuf[base+3];

    destbuf[base+3]=
          fv[ 3] * srcbuf[base+0]
        + fv[ 7] * srcbuf[base+1]
        + fv[11] * srcbuf[base+2]
        + fv[15] * srcbuf[base+3];
}

(4) は Android に含まれるライブラリを利用したものです。
最も低速でした。

// java : (4)
for( int i= 0 ; i< MAX_TRANSFORM ; i++ ){
    int base= i*4;
    Matrix.multiplyMV( destbuf, base, fv, 0, srcbuf, base );
}

●NDK

RenderScript (1) にかなり近い数値でした。
debug build ではこの半分の速度だったので release build の最適化が
かなり効いているようです。
展開コードを調べると、fmuls + fadds の組み合わせが最適化により
fmacs の積和命令に置き換わっていることがわかります。

// NDK C++ : (6)
static void loopndk2( const Vect4* vin, Vect4* vout, int length )
{
    for( int i= 0 ; i< length ; i++ ){
        vout->x= TransformMatrix._11 * vin->x
                +TransformMatrix._21 * vin->y
                +TransformMatrix._31 * vin->z
                +TransformMatrix._41 * vin->w;
        vout->y= TransformMatrix._12 * vin->x
                +TransformMatrix._22 * vin->y
                +TransformMatrix._32 * vin->z
                +TransformMatrix._42 * vin->w;
        vout->z= TransformMatrix._13 * vin->x
                +TransformMatrix._23 * vin->y
                +TransformMatrix._33 * vin->z
                +TransformMatrix._43 * vin->w;
        vout->w= TransformMatrix._14 * vin->x
                +TransformMatrix._24 * vin->y
                +TransformMatrix._34 * vin->z
                +TransformMatrix._44 * vin->w;
        vin++;
        vout++;
    }
}

試しにアセンブラで記述してみたのが (7)。
やはり NEON はなくても fmacs が効果的です。
さらに最速ケースを調べるためにスレッド化し、Cortex-A9 2 core を使って
並列実行したのが (8) です。

// NDK asm : (7)/(8)
    __asm__ __volatile( "\
    1: \n\
        vldmia %1!, {d8-d9} \n\
        vldmia %0, {d0-d7} \n\
        fmuls s20,s0 ,s16 \n\
        fmuls s21,s4 ,s16 \n\
        fmuls s22,s8 ,s16 \n\
        fmuls s23,s12,s16 \n\
        fmacs s20,s1 ,s17 \n\
        fmacs s21,s5 ,s17 \n\
        fmacs s22,s9 ,s17 \n\
        fmacs s23,s13,s17 \n\
        fmacs s20,s2 ,s18 \n\
        fmacs s21,s6 ,s18 \n\
        fmacs s22,s10,s18 \n\
        fmacs s23,s14,s18 \n\
        fmacs s20,s3 ,s19 \n\
        fmacs s21,s7 ,s19 \n\
        fmacs s22,s11,s19 \n\
        fmacs s23,s15,s19 \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"
    );

以上より

RenderScript(LLVM)
 ・Single Thread 実行
 ・rsMatrixMultiply() は内部で fmacs (積和) 命令を使用し最適化されている
 ・RenderScript は通常の演算は fmacs (積和) に展開しない

NDK(gcc)
 ・通常の演算コードも fmacs (積和) に最適化される

おそらく NEON があれば rsMatrixMultiply() が更に高速で、コンパイラが
展開したコードよりも高速に実行できたのではないかと考えられます。

●まとめ

SIMD がなくマルチコアも活用されないので、そのまま NDK C++ で書いた
コードと大差ない結果となってしまいました。
今後 Android OS 4.0 (Ice Cream Sandwich) が登場し、さらに多くの
デバイスで走るようになれば面白いデータが見られるのではないかと思います。

当然ながら、何でもできる NDK が最も速度を引き出せることには変わりありません。
その代わり全部自分でやらなければなりません。
armeabi, arameabi-v7a, neon, x86 等のアーキテクチャ対応、
SIMD 対応、スレッド化などすべてに取り組むのは大変です。

RenderScript はアーキテクチャを選ばず、バイナリ化してもランタイムの
改良で性能が向上する可能性があります。
これまで使ってきたように Java から気軽に呼び出せるため
Android 4.0 が普及すれば色々と使い道が増えると思います。

関連エントリ
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)

Android 3.x RenderScript (6) Graphics まとめ

Android には様々な描画手段が用意されています。
GPU を使って高速に描画可能な組み合わせは下記の通り。

View + onDraw (HW)
RSSurfaceView + RenderScript
GLSurfaceView + OpenGL ES 2.0 (Java)
GLSurfaceView + OpenGL ES 2.0 (NDK)
NativeActivity + OpenGL ES 2.0 (NDK)

下に行くほど低レベルな API になります。
実際に使ってみた限り、描画速度もほぼこの順番になるようです。

● RenderScript

マテリアルのセットアップやジオメトリ演算など、描画ループ自体を
Native 実行可能な C言語で記述することが可能です。
リソースの生成や初期化は Java で行い、プロジェクトも開発環境も
Java と非常に親和性が高いのが特徴です。

Compute 用の RenderScript を簡単に呼び出せます。
大量の並列演算を伴う場合も対応できます。

描画 API は OpenGL ES 2.0 の上に構築された Class ライブラリとなります。
理解してしまえば OpenGL ES 2.0 を直接扱うよりも簡単で使いやすいですが、
直接 GL 命令を使えないのでできることに限界があります。
特にぎりぎりまでハード性能を引き出すような用途には向きません。
また描画に GLSL を使う場合は OpenGL ES 2.0 の知識も必要となります。

利点
 ・C言語 + Native による高速実行。
 ・CPU アーキテクチャ非依存でどの環境でも走る。
 ・Java と親和性が高く Java のプロジェクトに組み込みやすい。
 ・描画 API が Class 化されており GL 直接より簡単。リソースも扱いやすい。
 ・rs_cl と併用できる。

欠点
 ・Android 3.0/API Level 11 以上のみ
 ・GL 周りで高速化に限界がある。
 ・描画リソースを一旦 script に渡してから描画する必要あり。

● GLSurfaceView + OpenGL ES 2.0 (Java)

GPU のハードウエアや OpenGL をよく知っていて、普段から GPU の性能を
めいっぱい使うような使い方をしている場合こちらの方が速いです。
CPU 負荷よりも GPU 負荷のほうが高いと考えられるためです。

逆にもし CPU がボトルネックとなっている場合は RenderScript の方が
向いているといえるでしょう。
アーキテクチャ非依存という点では同一で、特に OpenGL に慣れているなら
容易に扱えます。

利点
 ・CPU アーキテクチャ非依存でどの環境でも走る。
 ・Java だけで完結するため開発しやすい。
 ・GL API をそのまま使うため機能制限が無い

欠点
・Android 2.3/API Level 9 以上 (2.2 ではバグあり)
 ・低レベルな GL API をそのまま使うためコード量が増える
 ・Java によるボトルネックが考えられる

● GLSurfaceView + OpenGL ES 2.0 (NDK)

OpenGL ES 2.0 そのままで、かつ普通に C/C++ で開発できるため他の
プラットフォームで経験があるなら一番親和性が良いはずです。
ただし Java 環境への dll によるアドオンとなるため jni による通信が必要。
開発も両環境を行き来する必要が生じます。特にデバッグで。
また Native Code への事前コンパイルなので、ビルド時に実行環境が固定されます。
速度も一番引き出せる可能性がありますが、ハードに近いところまでそれなりの
知識が必要です。

利点
 ・C/C++ / OpenGL ES そのままなのでソース互換性が高い(移植しやすい,理解しやすい)
 ・GL API をそのまま使うため機能制限が無い
 ・Android 2.0/API Level 5 以上対応なので多くのデバイスで実行できる

欠点
 ・CPU アーキテクチャに依存する。ビルド時に実行環境が固定される
 ・低レベルな GL API をそのまま使うためコード量が増える
 ・リソースアクセスやイベント時に Java と通信が必要

● NativeActivity + OpenGL ES 2.0 (NDK)

Java を使わずに開発できます。
描画は EGL + OpenGL ES 2.0 のみで、全て自力で描画する必要があります。
Android の便利な Class ライブラリを使うことができません。
ほぼゲーム用。

利点
 ・C/C++ / OpenGL ES そのままなのでソース互換性が高い(移植しやすい,理解しやすい)
 ・C/C++ だけで完結し Java コードを必要としないためボトルネックがない。
 ・GL API をそのまま使うため機能制限が無い

欠点
・Android 2.3/API Level 9 以上のみ
 ・CPU アーキテクチャに依存する。ビルド時に実行環境が固定される
 ・低レベルな GL API をそのまま使うためコード量が増える
 ・使える API が限られており、Android の豊富なライブラリが使えない

続きます。次は RenderScript Compute の速度 「Android 3.x RenderScript (7) RenderScript Compute の速度」

関連エントリ
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)