16.78MHz -7ページ目
Amebaホームピグアメブロ
芸能人ブログ人気ブログ
新規登録
ログイン

ブログ画像一覧を見る

このブログをフォローする

<< 前のページへ最新 | 5 | 6 | 7 | 8 | 9 | 最初次のページへ >>

nvcc特有のデバッグに関する問題

CUDAのコンパイラnvccはデバイス側のコードをコンパイルする時グラフィックハードウェアのシンプルさを最大限に活かした大胆な最適化を行う。このことがバグを探すときに思わぬ落とし穴になることがある。
以下のようなコードを考えてみる。
extern "C"
  __global__
  void entry( float *_output, float *_input ) {
    float hoge = _input[ threadIdx.x - 1 ];
    _output[ threadIdx.x ] = hoge + 1;
  }

今_inputには素直にcuMemAllocで確保したグローバルメモリの先頭アドレスが渡されているとすると、このコードには明らかなバグがある。threadIdx.xが0になるスレッドでメモリの範囲外参照が起こるからだ。
新しめのCUDAなら範囲外参照が起こるとlaunch failedエラーで__global__関数の実行を停止させるのでバグの存在はすぐに判明する。次は具体的にどこに問題があるのかを探すことになるわけだが、こういうとき便利なデバイスエミュレーションモードはドライバAPI+cubinでCUDAをいじる前提で開発している場合使用できない。そういった場合この例のようなシンプルなコードならばともかく、それなりの大きさがある場合はとりあえず問題箇所を絞り込むために、後ろの方をコメントアウトして最初の方だけ実行してみることは適切なように思える。

extern "C"
  __global__
  void entry( float *_output, float *_input ) {
    float hoge = _input[ threadIdx.x - 1 ];
//    _output[ threadIdx.x ] = hoge + 1;
  }

このコードで問題があるのはコメントアウトされた行の前の行なので、これをコンパイルして実行すると先ほどと同じく範囲外参照が起こるはずだが、実際にはこのコードは実行に成功する。問題のある箇所が特定出来ていない状態でこんな結果が出たら、開発者は問題箇所の絞り込みを誤ってしまうだろう。なぜこのコードの実行が成功するのかはnvccが出力したアセンブラコードを見ればはっきりする。
$ nvcc --ptx hoge.cu
hoge.cu(4): warning: variable "hoge" was declared but never referenced
$ cat hoge.ptx
...
{
.loc 15 3 0
$LBB1_entry:
.loc 15 6 0
exit; //
$LDWend_entry:
} // entry

速攻exit

つまり、範囲外参照して初期値を決定した変数hogeがその後メモリ上のどの値にも全く影響を与えていないために、この部分の処理はまるごと削除されてしまい、結果として起こるはずの範囲外参照が起こらなかったのだ。
CUDAのデバイス側コードではこのようにデバッグ目的でソースの一部だけを実行しようとすると、実際には実行されてほしい部分が実行されないことがある点に注意が必要だ。

Router Station Pro

結構昔からiptablesが本家カーネルにマージされていたこともあり、Linux PCをルータ化するというのはその手の雑誌の定番のネタだった。自分で作る高機能ルータは確かに魅力満点なのだが、普通のPCというものはえてしてルータとして使うには、大きすぎて、うるさすぎて、消費電力が大きすぎるものだった。近年はAtomプロセッサを搭載した省電力マザーボードのような製品も増えてきたため、そうした製品を使うのも一つの手ではあるのだが、世の中にはLinuxルータ専用ボードなんてものが存在する。

Router Station ProはUbiquitiが開発した組み込みLinuxボードで、ルータにせずして何に使うんだと言わんばかりのコネクタが特徴。具体的なスペックは以下のとおり

680MHz MIPSプロセッサ
メインメモリ128MB
フラッシュメモリ16MB
SDカードスロット
ギガビットイーサネット x2(うち一方はハブで3ポート)
USB 2.0 x1

イマドキのシングルボードコンピュータにありがちな派手な映像出力等は備えていないが内蔵フラッシュメモリにOpen-WRTが書き込まれており、買ったその日からLinuxルータとして使うことが出来る。拡張端子にはオプションで無線LAN等を追加することが出来るがおそらく日本では使用できない。ただ、この拡張端子はただのmini-PCIのようなので、他のデバイスを増設するために使える可能性はある。

Ubiquitiは直販はやっていないので、こことかこことかここあるいはこんなとこから買うことになる。どれも本家サイトに書かれている価格より地味に高いが気にしない。

SAM9G45

組み込みLinuxナードにはお馴染みのLinux for DevicesによるとAtmelから$8のARMプロセッサが登場するらしい。
価格からしてマイコンやPICを置き換えるような類の製品だろうと思ったのだが、

ブロック図

こいつ、できる！

さて、驚きのプロセッサが登場したわけだが、BGAのチップ単品で売られても個人ではどうすることも出来ない。BeagleBoardのような安価なボードが存在しないものかと思い調べたところ、あった。

簡単にスペックをまとめると
ARM926EJ-S 400MHz
128MB DDR2-SDRAM
256MB フラッシュメモリ
USB 2.0 x4
100BASE-TX Ethernet
SDカード端子
480x272(PSP液晶サイズ)液晶コネクタ
大きさがPico-ITXマザーボード

と、何かと人気のOMAP3ベースのボードと比較するといまいちパッとしない性能なものの、$69と安価な点と、最初からAndroidがインストールされていて組み込みLinux初心者でも取り合えず起動で挫折することはない点がポイント。

リリースは今年の年末の見通し。

メイン変数

#ifdef __i386__
const int main = 0x909090c3; // ret
#endif
#ifdef __PPC__
const int main = 0x4e800020; // blr
#endif
#ifdef __ARMEL__
const int main = 0xe12fff1e; // bx     lr
#endif

$ gcc main_the_variable.c -o main_the_variable -Wall
main_the_variable.c:2: 警告: ‘main’ is usually a function
$ ./main_the_variable
$

普通に動くから困る。
そもそも動く環境が限られているからintの長さが4バイトじゃなかったら、とかそういうことは考えていない。

gccの最適化オプションでVFP/NEONを使う

古くからi387のようなハードウェアがあったx86とは異なりARMには最近までハードウェアで浮動小数点数を計算する仕組みがなかった。
このためARM用のgccはデフォルトで浮動小数点数の演算をソフトウェアでエミュレーションするバイナリを吐く。
かつてXScaleを積んだシャープのZaurus(通称リナザウ)にスケーラブルフォントやOgg Vorbisといった浮動小数演算を伴う処理を多用するKDE3.4を入れてみたところ動いたが実用に耐えないほど遅かった、という話があったが原因のひとつがこの浮動小数点数のエミュレーションにあったのではないかと思う。
あれから4年

今日のハイエンドなARMプロセッサにはVFP(ARMv6)あるいはNEON(ARMv7、VFP上位互換)といった浮動小数点数SIMD命令セットが搭載されており、浮動小数点数をハードウェアで演算することができる。
BeagleBoardはじめとするTI OMAP3搭載マシンはNEONに対応しているため、これを使わない手は無い。
さっそくその威力を見てみようと思いサンプルプログラムを用意した。
#include <stdio.h>
int main(){
int counter0, counter1;
double value;
  for( counter0 = 0; counter0 != 1000; counter0++ )
    for( counter1 = 0; counter1 != 1000; counter1++ )
      value++;
  printf( "%lf", value );
  return 0;
}

コンパイルして実行時間を計測
$ gcc test_fp.c -o disable_neon
$ time ./disable_neon
1000000.000000
real    0m0.112s
user    0m0.109s
sys     0m0.000s
$ gcc -march=armv7-a -mtune=cortex-a8 -mfpu=neon \
      test_fp.c -o enable_neon
$ time ./enable_neon
1000000.000000
real    0m0.112s
user    0m0.102s
sys     0m0.008s

...あれ?　同じ?
何か予想と違うことが起こってそうだったのでコンパイラが吐いたアセンブラコードを見てみた。
$ gcc test_fp.c -o disable_neon.s -S
$ cat disable_neon.s
...
.L4:                                               
        ldrd    r0, [fp, #-20]                     
        mov     r2, #0                             
        mov     r3, #1069547520                    
        add     r3, r3, #3145728                   
        bl      __aeabi_dadd                       
        mov     r2, r0                             
        mov     r3, r1                             
        strd    r2, [fp, #-20]                     
        ldr     r3, [fp, #-24]                     
        add     r3, r3, #1                         
        str     r3, [fp, #-24]
...
$ gcc -march=armv7-a -mtune=cortex-a8 -mfpu=neon \
      test_fp.c -o enable_neon.s -S
$ cat enable_neon.s
...
.L4:                                               
        ldrd    r0, [fp, #-20]                     
        mov     r2, #0                             
        mov     r3, #1069547520                    
        add     r3, r3, #3145728                   
        bl      __aeabi_dadd                       
        mov     r2, r0                             
        mov     r3, r1                             
        strd    r2, [fp, #-20]                     
        ldr     r3, [fp, #-24]                     
        add     r3, r3, #1                         
        str     r3, [fp, #-24]
...
$ diff disable_neon.s enable_neon.s
$

gcc-4.3はNEONに対応しているはずなのだが、何故かNEONを使用するように指定しても全く同じソフトウェアによる浮動小数点数演算を行うコードを吐いていた。
色々調べたところ
NEON fpu support by setting -mfpu=neon -mfloat-abi=softfp

という話を見つけたので試しに-mfloat-abi=softfpを付けてみた。
$ gcc -march=armv7-a -mtune=cortex-a8 -mfpu=neon \
      -mfloat-abi=softfp \
      test_fp.c -o enable_neon.s -S
$ cat enable_neon.s
...
.L4:                                               
        fldd    d17, [fp, #-20]                    
        fconstd d16, #112                          
        faddd   d16, d17, d16                      
        fstd    d16, [fp, #-20]                    
        ldr     r3, [fp, #-24]                     
        add     r3, r3, #1                         
        str     r3, [fp, #-24]
...
$ gcc -march=armv7-a -mtune=cortex-a8 -mfpu=neon \
      -mfloat-abi=softfp \
      test_fp.c -o enable_neon
$ time ./enable_neon
1000000.000000
real    0m0.092s
user    0m0.086s
sys     0m0.000s

おぉ、使ってる。
最適化をかけると差はもっと顕著に出てくる。

$ gcc -Os -mabi=aapcs-linux -fforce-addr \
      -fomit-frame-pointer -fstrength-reduce \
      -fno-strict-aliasing \
      test_fp.c -o disable_neon
$ time ./disable_neon
1000000.000000
real    0m0.096s
user    0m0.094s
sys     0m0.000s
$ gcc -Os -march=armv7-a -mtune=cortex-a8 -mfpu=neon \
      -mfloat-abi=softfp -mabi=aapcs-linux -fforce-addr \
      -fomit-frame-pointer -fstrength-reduce \
      test_fp.c -o enable_neon
$ time ./enable_neon
1000000.000000
real    0m0.025s
user    0m0.023s
sys     0m0.000s
<< 前のページへ最新 | 5 | 6 | 7 | 8 | 9 | 最初次のページへ >>

ブログ画像一覧を見る

このブログをフォローする