Cortex-A53
(3ヶ月くらいまえに)携帯あたらしくしてCortex-A53(ただし KitKat なので32bit)にしたのだけど、全然性能はかってないのでなんかはかろうと思って
https://github.com/tanakamura/instruction-bench
のARM 版をつくった
https://github.com/tanakamura/ude-bench
今はまだコマンドラインでしか動かないので、Makefileなんかうまいこと修正して、できたinstbench を adb shell とかで動かしてなんとかしておいて。UIはそのうちなんとかすると思う…。
あと Raspi2(Cortex A7) 持ってて動かしてみたという人がいたら、結果なんらかの方法で送ってもらえると嬉しいです(もしかしたらperf_event_open無くて動かないかも)。
よく考えたら解説書いたことなかったので書いておくと、
loop: add r0, r0, r0 add r0, r0, r0 add r0, r0, r0 ... add r0, r0, r0 sub r10, 1 bne loop
みたいなのを動かして、命令のレイテンシを、
loop: add r0, r0, r0 add r0, r1, r1 add r0, r2, r2 ... add r0, r7, r7 sub r10, 1 bne loop
みたいなのを動かして、命令のスループットを計測する。
CPI、 IPC は同じ値を示してて、単なる逆数。直感的には、latency のときはCPI、throughput のときはIPCを見たほうが理解しやすいと思う。サイクル値は、パフォーマンスカウンタから読んでるので、省電力機能の影響は受けてないはず。
Cortex A53 だと、↓ だった。
== num_insn = 256 ==
generic : add rd, rm, rn : latency : CPI= 1.00, IPC= 1.00
generic : add rd, rm, rn : throughput : CPI= 0.51, IPC= 1.96
generic : add rd, rm, rn : rename : CPI= 0.51, IPC= 1.97
generic : adds rd, rm, rn : latency : CPI= 1.00, IPC= 1.00
generic : adds rd, rm, rn : throughput : CPI= 0.51, IPC= 1.97
generic : adds rd, rm, rn : rename : CPI= 0.51, IPC= 1.96
generic : add rd, rm, rn, lsl #4 : latency : CPI= 2.00, IPC= 0.50
generic : add rd, rm, rn, lsl #4 : throughput : CPI= 0.51, IPC= 1.97
generic : add rd, rm, rn, lsl #4 : rename : CPI= 0.51, IPC= 1.96
generic : add rd, rm, imm : latency : CPI= 1.00, IPC= 1.00
generic : add rd, rm, imm : throughput : CPI= 0.51, IPC= 1.97
generic : add rd, rm, imm : rename : CPI= 0.51, IPC= 1.97
generic : add rd, rm, pc : latency : CPI= 1.01, IPC= 0.99
generic : add rd, rm, pc : throughput : CPI= 0.51, IPC= 1.96
generic : add rd, rm, pc : rename : CPI= 0.51, IPC= 1.97
generic : add pc, pc, 0 : latency : CPI= 4.52, IPC= 0.22
generic : orr rd, rm, rn : latency : CPI= 1.00, IPC= 1.00
generic : orr rd, rm, rn : throughput : CPI= 0.51, IPC= 1.97
generic : orr rd, rm, rn : rename : CPI= 0.51, IPC= 1.97
generic : eor rd, rm, rn : latency : CPI= 1.00, IPC= 1.00
generic : eor rd, rm, rn : throughput : CPI= 0.51, IPC= 1.97
generic : eor rd, rm, rn : rename : CPI= 0.51, IPC= 1.97
generic : mul rd, rm, rs : latency : CPI= 3.00, IPC= 0.33
generic : mul rd, rm, rs : throughput : CPI= 1.00, IPC= 1.00
generic : mul rd, rm, rs : rename : CPI= 1.00, IPC= 1.00
generic : mla rd, rm, rs, rn : latency : CPI= 3.00, IPC= 0.33
generic : mla rd, rm, rs, rn : throughput : CPI= 1.00, IPC= 1.00
generic : mla rd, rm, rs, rn : rename : CPI= 1.00, IPC= 1.00
generic : ldr rt, [rn, rm] : latency : CPI= 3.00, IPC= 0.33
generic : ldr rt, [rn, rm] : throughput : CPI= 1.00, IPC= 1.00
generic : ldr rt, [rn, rm] : rename : CPI= 1.00, IPC= 1.00
generic : ldr rt, [rn, rm, lsl #4] : latency : CPI= 5.01, IPC= 0.20
generic : ldr rt, [rn, rm, lsl #4] : throughput : CPI= 3.01, IPC= 0.33
generic : ldr rt, [rn, rm, lsl #4] : rename : CPI= 3.01, IPC= 0.33
generic : ldm rt, {r0} : throughput : CPI= 1.01, IPC= 0.99
generic : ldm rt, {r0} : rename : CPI= 1.01, IPC= 0.99
generic : ldm rt, {r0-r3} : throughput : CPI= 2.01, IPC= 0.50
generic : ldm rt, {r0-r3} : rename : CPI= 2.01, IPC= 0.50
generic : ldm rt, {r0-r7} : throughput : CPI= 4.01, IPC= 0.25
generic : ldm rt, {r0-r7} : rename : CPI= 4.01, IPC= 0.25
generic : ldrex rd, [rn] : throughput : CPI= 3.00, IPC= 0.33
generic : ldrex rd, [rn] : rename : CPI= 3.00, IPC= 0.33
generic : strex rd, rm, [rn] : throughput : CPI= 3.00, IPC= 0.33
generic : strex rd, rm, [rn] : rename : CPI= 3.00, IPC= 0.33
generic : ldrex r0, [rn]; strex rd, r0, [rn] : throughput : CPI= 12.02, IPC= 0.08
generic : ldrex r0, [rn]; strex rd, r0, [rn] : rename : CPI= 12.02, IPC= 0.08
generic : str rt, [rn, #0] : throughput : CPI= 1.00, IPC= 1.00
generic : str rt, [rn, #0] : rename : CPI= 1.00, IPC= 1.00
generic : {str->ldr}->... : latency : CPI= 3.00, IPC= 0.33
generic : {strb->ldr}->... : latency : CPI= 3.00, IPC= 0.33
neon64 : vadd.f32 d, d, d : latency : CPI= 4.01, IPC= 0.25
neon64 : vadd.f32 d, d, d : throughput : CPI= 1.00, IPC= 1.00
neon64 : vadd.f32 d, d, d : rename : CPI= 0.51, IPC= 1.97
neon128 : vadd.f32 q, q, q : latency : CPI= 4.01, IPC= 0.25
neon128 : vadd.f32 q, q, q : throughput : CPI= 1.00, IPC= 1.00
neon128 : vadd.f32 q, q, q : rename : CPI= 1.01, IPC= 0.99
neon64 : vmul.f32 d, d, d : latency : CPI= 4.01, IPC= 0.25
neon64 : vmul.f32 d, d, d : throughput : CPI= 1.00, IPC= 1.00
neon64 : vmul.f32 d, d, d : rename : CPI= 0.51, IPC= 1.96
neon128 : vmul.f32 q, q, q : latency : CPI= 4.01, IPC= 0.25
neon128 : vmul.f32 q, q, q : throughput : CPI= 1.01, IPC= 0.99
neon128 : vmul.f32 q, q, q : rename : CPI= 1.01, IPC= 0.99
neon64 : vmul.f32 d, d, d : latency : CPI= 4.01, IPC= 0.25
neon64 : vmul.f32 d, d, d : throughput : CPI= 1.00, IPC= 1.00
neon64 : vmul.f32 d, d, d : rename : CPI= 0.51, IPC= 1.97
neon128 : vmul.f32 q, q, q : latency : CPI= 4.01, IPC= 0.25
neon128 : vmul.f32 q, q, q : throughput : CPI= 1.01, IPC= 0.99
neon128 : vmul.f32 q, q, q : rename : CPI= 1.01, IPC= 0.99
neon64 : vmla.f32 d, d, d : latency : CPI= 4.01, IPC= 0.25
neon64 : vmla.f32 d, d, d : throughput : CPI= 1.00, IPC= 1.00
neon64 : vmla.f32 d, d, d : rename : CPI= 0.51, IPC= 1.97
neon128 : vmla.f32 q, q, q : latency : CPI= 4.01, IPC= 0.25
neon128 : vmla.f32 q, q, q : throughput : CPI= 1.01, IPC= 0.99
neon128 : vmla.f32 q, q, q : rename : CPI= 1.01, IPC= 0.99
neon64 : vld1.32 d, [rn] : throughput : CPI= 1.00, IPC= 1.00
neon64 : vld1.32 d, [rn] : rename : CPI= 1.01, IPC= 0.99
neon64 : vld2.32 d, [rn] : throughput : CPI= 1.01, IPC= 0.99
neon64 : vld2.32 d, [rn] : rename : CPI= 1.01, IPC= 0.99
neon128 : vld4.32 q, [rn] : throughput : CPI= 2.01, IPC= 0.50
neon128 : vld4.32 q, [rn] : rename : CPI= 2.01, IPC= 0.50
neon64 : vst1.32 d, [rn] : throughput : CPI= 1.00, IPC= 1.00
neon64 : vst1.32 d, [rn] : rename : CPI= 1.01, IPC= 0.99
neon64 : vst2.32 d, [rn] : throughput : CPI= 1.01, IPC= 0.99
neon64 : vst2.32 d, [rn] : rename : CPI= 1.01, IPC= 0.99
neon128 : vst4.32 q, [rn] : throughput : CPI= 2.01, IPC= 0.50
neon128 : vst4.32 q, [rn] : rename : CPI= 2.01, IPC= 0.50
neon64 : vcvt.f32.s32 d, d : throughput : CPI= 0.51, IPC= 1.97
neon64 : vcvt.f32.s32 d, d : rename : CPI= 0.51, IPC= 1.97
neon128 : vcvt.f32.s32 q, q : throughput : CPI= 1.01, IPC= 0.99
neon128 : vcvt.f32.s32 q, q : rename : CPI= 1.01, IPC= 0.99
neon64 : vcvt.s32.f32 d, d : throughput : CPI= 0.51, IPC= 1.97
neon64 : vcvt.s32.f32 d, d : rename : CPI= 0.51, IPC= 1.97
neon128 : vcvt.s32.f32 q, q : throughput : CPI= 1.01, IPC= 0.99
neon128 : vcvt.s32.f32 q, q : rename : CPI= 1.01, IPC= 0.99結構普通だなーという感じだけど(まあ今時のCPUなら普通なのが当然だが…)、いくらか気になる点は、
オペランドのレジスタをシフトするとおそい。特にアドレスをシフトするとスループット1/3 になる。
vcvt.f32.s32 はレイテンシ1だがなんかの間違いな気がするがこんなもんだっけ…
throughput と rename のちがいは、
// throughput loop: vadd.f32 d0, d1, d1 vadd.f32 d0, d2, d2 ... vadd.f32 d0, d8, d8 sub r10, 1 bne loop
// rename loop: vadd.f32 d1, d1, d1 vadd.f32 d2, d2, d2 ... vadd.f32 d8, d8, d8 sub r10, 1 bne loop
なのだが、64bit neon だと、なぜかこれで性能が変わる。rename のほうはスループット2出るが、throughputのほうは出ない。
vmla.f32 q,q,q はスループット1出るが、vld.f32 q, [r] はスループット0.5。
あたりかな…
あと recp, rsqrt, zip, シフト全般などを追加したほうがいい気がする。(追加するの結構ダルいのでやらないかもしれない)
clminibench for Android
上のをUI作るまえに復習しようと思って、clminibench をちゃんとリリースしておいた。
https://play.google.com/store/apps/details?id=main.jp.Int.clminibench
作ったとき一緒に作った資料↓。
http://int.main.jp/txt/mali-chang.pdf
もう二年前の話か…ていうか、この時の記録を書いてない気がした。
主に、clEnqueueNDRangeKernel のレイテンシが 10usec なのか、100usec なのか知りたい人向けのベンチマークなので、あまり期待しないでね…
Androidは OpenCL icd に対応してなくて、libOpenCL.so の位置を自分で探さないといけないのだが、どこにも資料が無くて、一応、
- /system/vendor/lib/libOpenCL.so
- /system/vendor/lib/egl/libGLES_mali.so
のふたつはロードを試みてるが、それ以外に置いてあると動かない。
あと、int.main.jp を package 名にすると、int がキーワードだからビルドできないという悲しい話があった。
関東GPGPU勉強会 #2
というわけで、せっかくだから記録を書いておくと…覚えてるわけが無いのだった。(http://togetter.com/li/511911 多分こっちを見たほうが情報量がある)
このときは、Mali-T604 は19.0GFLOPSと書いてるのだけど、どこで見たか忘れたけど、72か74ぐらいが理論値らしい。ただ、この単純にmad4を並べた書きかたで出ないなら、どうやったら出るんだという気がする。
気になる記述としては、
に、「sinとかcosに含まれる演算を計算性能に含めるのはセコいだろ(まあでもdot productのは含めてもいいかな…)(超意訳)」とか書いてあって、なんかdot() なら性能出るのかも、という気がした。たぶん、パイプライン上に複数の演算があって、特定の書きかたしたら、複数の演算が1cycleでできるようになってるとかなのではないかな…
しかしNexus10 のOpenCLはつぶされたのでもう確認できない。
