はじめまして、ネットエージェント株式会社、取締役の愛甲健二と申します。今後、このブログにて主に技術的な内容について書かせていただきたいと考えていますので、ぜひよろしくお願いいたします。
今日はGPGPUについて少し紹介させていただきたく思います。
-----
GPU（Graphics Processing Unit）は画像処理を専門とする演算装置ですが、この演算資源をより汎用的なプログラム処理に応用しようという技術のことをGPGPU（General-purpose computing on GPUs）と呼びます。
GPUはCPUとは異なる命令セットを持つため、CPUと同じ感覚でGPUへ処理を渡すことは難しいですが、近年はCUDAを初めとするGPGPU用のライブラリがいくつか提供されており、より手軽にGPUを利用できます。
CUDAを用いて作られた実行ファイルは、内部にGPU用の命令コード（正確にはCUDA用の命令コード:PTX）を持ちます（右図参照）。
この命令コードはCUDAライブラリ内部で各GPUの命令コードへ変換され、GPU用のメモリ空間へコピーされ、CPUからGPUへ処理の開始通知が送られてからGPUが命令コードを実行します。実行結果は再びGPU用のメモリから普通のメモリ（主記憶）へコピーされます。
CUDAの場合は、GPUリソースへのアクセスはすべてDLL（cutil32.dll等のCUDAライブラリ）が請け負ってくれますので、関数を呼び出すだけでGPUを利用できます。
GPU側の処理は、デバッガを使って動的に実行状態を追いかけることをはじめ、命令コード自体の解析が困難という特徴があるため、新しい難読化の手法としてGPUを用いた方法が出現しています。具体的には、2008年にDaniel Reynaud氏がセキュリティカンファレンスruxcon 2008にて「GPU Powered Malware」というプレゼンテーションを発表しました。これは、まだ有用なGPU用のデバッガやエミュレータが開発されていない現状においては、GPUを使ったソフトウェアの解析は非常に困難であり、それがマルウェアの進化に影響を及ぼす、という内容でした。
もちろん、2010年の現在においても、GPU用命令コードの解析環境が十分に整っているとは言い難く、マルウェアを解析する側としても辛いところですが、ことCUDAに関してはdecudaと呼ばれる逆アセンブラがすでに存在します。
decudaはCUDA用のアセンブリ言語であるPTXを出力します。試しにCUDAで作成した実行ファイルを逆アセンブルしてみます。
$ python decuda.py TEST.cubin
000000: 1000c801 0423c780 mov.b32 $r0, s[0x0010]
000008: d00e000d 80c00780 mov.u32 $r3, g[$r0]
000010: 100d8009 0ccccccf mov.b32 $r2, 0xcccccccd
000018: a0000605 04114780 cvt.rn.abs.u32.s32 $r1, $r3
000020: 40050411 00000780 mul24.lo.u32.u16.u16 $r4, $r1.lo, $r2.hi
000028: 60040611 000107c0 mad24.lo.u32.u16.u16.u32 $p0|
$r4, $r1.hi, $r2.lo, $r4
000030: 30100815 c4100780 shl.u32 $r5, $r4, 0x00000010
000038: 30100811 e4100780 shr.u32 $r4, $r4, 0x00000010
000040: 600405fd 000147d8 mad24.lo.u32.u16.u16.u32 $p1|
$o127, $r1.lo, $r2.lo, $r5
000048: 21000811 04400880 @$p0.cf add.u32 $r4, $r4, c1[0x0000]
000050: 60050605 0c011780 mad24c1.lo.u32.u16.u16.u32 $r1,
-$r1.hi, $r2.hi, -$r4
000058: 30020205 ec100780 shr.s32 $r1, $r1, 0x00000002
000060: 307c07fd 6c0047c8 set.lt.s32 $p0|$o127, $r3, $r124
000068: 30008209 00000003 subr.b32 $r2, $r1, 0x00000000
000070: 10000405 0403c280 @$p0.ne mov.b32 $r1, $r2
000078: d00e0005 a0c00781 mov.end.u32 g[$r0], $r1
見ての通りx86系CPUとは異なります。上記のコードは、以下のC言語のプログラムをコンパイルして逆アセンブルした結果です。
__global__ void cudaKernel(int *gpu)
{
gpu[0] /= 0x05;
}
また最近になって、NVIDIAは「NVIDIA Parallel Nsight」と呼ばれるGPU用デバッガを開発、リリースしています。これはVisual Studio上でCUDAのカーネル（GPU用コード）を動的デバッグできるというもので、これを使えるならばよりGPUの解析が容易になると思います。
GPUを用いる場合は、どこかに必ずGPU用の命令コードを持たなければならないため、解析不可能という状態にはならないと思いますが、ただCPUとは様々な部分で仕様が異なりますので、実際問題、GPUを本格的に使われると解析は難しくなります。
また、GPUは専用のメモリ空間を持つので、データを隠す場所という意味でも利用できるかもしれません。メモリフォレンジックの観点になりますが、GPU用メモリは通常の方法ではメモリダンプできないため、この点においても今後リサーチが進むかもしれません。
このように、GPUを利用したソフトウェアの解析技術については、まだまだ発展途上であるため、今後も継続して調査研究を続けていきたいと思います。
