c-bata web

@c_bata_ のメモ。python と Go が多めです。

Wrapping up 2017 and 2018, Hello 2019.

そういえば振り返り記事書いてなかったなと思いたち書きました。 よくみたら昨年も忙しすぎて書いてなかった。

(写真は年始に会社の人と新潟まで泊まりでスノボいったときの景色)

Wrapping up 2017 and 2018

2018は振り返りをしてなかったので、2年前の目標を確認。

2017年抱負(エンジニアリング) - c-bata web

OSS

2017と2018に公開したOSS(example系は除く)を並べるとこんな感じ。もっと色々書いてたきがしますが公開まで持っていけなかったものも多いです。

Goの勉強会であまり登壇してないので全然知られてないと思いますが、仕事でGoを書いている影響かOSSPythonよりもGoのソフトウェアばかり書いてました。 go-prompt が2000stars近くまで伸び、RancherやKubicorn、Apache CloudMonkeyなど大手のOSSでも採用事例が出てきました。 kube-promptもGotime.fmDocker/Kubernetes 実践コンテナ開発入門 で紹介されたりと、かなりちゃんと使われるソフトウェアになってきたのかなと思います。

登壇

執筆にかなりの時間を使っていたわりには、意外と数こなしていました。 大きめのイベントでロングトークをしたのは、PyCon JP 2018、Django Congress JP 2018、AbemaTV Developer Conferenceの3つぐらいです。

あとPyCon JP 2016のときに用意した↓の資料が今年の8月頃かなり読まれるようになり、400はてぶくらいついたのでいろいろアップデートしました。 手前味噌ですがWSGIフレームワークの解説としては、英語圏の記事と比べてもわりとよくかけている気がしているので興味ある方はぜひ

Hello 2019!

4-5年後にはデータ系にキャリアをよせて行こうと思っていたのですが、社内のハイパーパラメーター最適化の研究をしているチームから声をかけてもらい面白そうだったので異動することにしました。3人しかいない今のAbemaTV配信チームから異動することを許可してくれたリーダー陣に感謝してます。データ系出身でもない学士の自分に声がかかる機会はもう今後そんなにないだろうということで異動申請を出すことにしました。

あと社内ISUCONの問題作成・ベンチマーカーの開発を仕切ることになりました。異動もあるのでメインの業務と両立できるか少し不安もありますが、楽しみながらこなしていきたい。今年の目標は次のあたり。

  1. 執筆のしごとをやりきる.
  2. 異動先で結果を出す
  3. 社内ISUCONを無事に成功させる

異動が正式に決まれば今後はOSSの開発に業務で取り組むことが多くなるので、仕事内容が趣味の開発にかなり近くなるような気がしています。 知り合いから(あと社内の人からも)なんで今の会社にしたのかわからないって結構言われるんですが、ミスったとか後悔してるって思いはいまのところ全くなくてまだ当分は転職せずに残ろうかなという気持ちです。

CUDAメモ: NVIDIA Tesla K80, CUDA 10.0, Ubuntu 18.04

たまたま参加したもくもく会に石本さんも来ていたので、疑問点があれば聞けるいい機会かなと思いCUDAの勉強をしてみました。 公式ドキュメント が非常に丁寧に書かれていてすごくスムーズに調査できた。最適化や行列計算のアルゴリズム的なテクニックに関しては、石本さんにおすすめされた CUDA C プロフェッショナル プログラミング (impress top gear) を今度読んで見る。

CUDA Toolkitのインストール

GCPNVIDIA Tesla K80のGPUをつけたインスタンスでいろいろ触ってみました。 インストール前にいくつか確認。 ubuntuのバージョン, cpu architecture, gccのバージョンがサポート対象であることを確認する。

$ lspci | grep -i nvidia
00:04.0 3D controller: NVIDIA Corporation GK210GL [Tesla K80] (rev a1)
$ uname -m && cat /etc/*release | head -n 2
x86_64
DISTRIB_ID=Ubuntu
DISTRIB_RELEASE=18.04
$ sudo apt-get install build-essential
$ gcc --version
gcc (Ubuntu 7.3.0-27ubuntu1~18.04) 7.3.0
$ uname -r
4.15.0-1023-gcp
$ sudo apt-get install linux-headers-$(uname -r)
…
linux-headers-4.15.0-1023-gcp is already the newest version (4.15.0-1023.24).

まとめると今回の環境は次の通り。

  1. ubuntu18.04
  2. gcc 7.3.0
  3. glibc 2.27
  4. linux kernel headers 4.15.0

CUDA 10.0は、今回使用している環境をサポートしている。次はいよいよCUDA Toolkitのインストールを行う。 公式サイトによるとdistribution-independent package (runfile packages)とdistribution-specific packagesの2つの方法があり後者がおすすめらしい。Ubuntuなのでcuda-downloadsのページから次の設定でdevパッケージを選択する。

f:id:nwpct1:20181110041229p:plain

See https://developer.nvidia.com/cuda-downloads?target_os=Linux&target_arch=x86_64&target_distro=Ubuntu&target_version=1804&target_type=debnetwork

$ wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/x86_64/cuda-repo-ubuntu1804_10.0.130-1_amd64.deb
$ sudo dpkg -i cuda-repo-ubuntu1804_10.0.130-1_amd64.deb
$ sudo apt-key adv --fetch-keys https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/x86_64/7fa2af80.pub
$ sudo apt-get update
$ sudo apt-get install cuda
...
done
$ nvcc -V
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2018 NVIDIA Corporation
Built on Sat_Aug_25_21:08:01_CDT_2018
Cuda compilation tools, release 10.0, V10.0.130

インストールも特に詰まることなく成功した。 一応 Installation Guide Linux :: CUDA Toolkit Documentation に沿って確認していく。

$ ls /usr/local/cuda-10.0/bin
bin2c        crt       cuda-gdbserver                cuda-memcheck  cuobjdump  gpu-library-advisor  nsight_ee_plugins_manage.sh  nvcc.profile  nvlink  nvprune  ptxas
computeprof  cuda-gdb  cuda-install-samples-10.0.sh  cudafe++       fatbinary  nsight               nvcc                         nvdisasm      nvprof  nvvp
$ cat .bash_profile
export PATH=/usr/local/cuda-10.0/bin${PATH:+:${PATH}}
export LD_LIBRARY_PATH=/usr/local/cuda-10.0/lib64${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}
$ source .bash_profile

CUDA 10.0のsamplesがあるので、それを展開していく。 CUDA Toolkitのversionに一致するsample集が簡単に手に入るのありがたい。

$ cuda-install-samples-10.0.sh .
Copying samples to ./NVIDIA_CUDA-10.0_Samples now...
Finished copying samples.
$ cd ./NVIDIA_CUDA-10.0_Samples
$ make

CPUケチって1 coreにしたんですが、全部やると15分ぐらいかかったので1つずつコンパイルしたほうがよかった。 次立てるときはcpu増やそう...

また後々NVIDIA Visual Profilerを見るために手元のmacOSにもCUDA Toolkitだけいれた。 CUDA Toolkit 10.0 Download | NVIDIA Developer からmac用のインストーラーをダウンロードして起動する。 「CUDA Driver」「CUDA Toolkit」「CUDA Samples」を選択できるが、NVIDIAGPUはないのでDriverは入れずにCUDA Toolkitだけ選択してインストールした。

CUDA Cプログラミング

よさそうなsampleとCUDA Toolkit Documentation のProgramming guideから読みはじめる。 ベクトル演算系がいいかなと思い、サンプル内の 0_Simple/vectorAdd から読み始めた。 コードを読み始めて気になったのは、 __global____device____host__ など見慣れない関数があったりdevice(gpu)上で実行される関数の呼び出しが特徴的な点があったりはするが、基本的なベクトルの加算処理はそれほど難しいことはなかった。基本的な流れは次の通り。

  1. Declare and allocate host and device memory.
  2. Initialize host data.
  3. Transfer data from the host to the device.
  4. Execute one or more kernels.
  5. Transfer results from the device to the host.

An Easy Introduction to CUDA C and C++

cudaMallocでdevice上にメモリー領域をアロケートし、cudaMemcpyでhostとdevice(gpu)間のデータ転送を行う。 ブロックサイズとブロック数を決めて呼び出しをして、結果をホストに転送する。 大きな行列(配列)を一度に演算することはできないので、呼び出しの際にどれくらいの単位で何個に分割するかを指定する。 func<<<N, M>>>(args...) という形式みたい (ex: vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);)。 詳細は Programming Guide :: CUDA Toolkit Documentation にかかれている。

このプログラムではfloat型をcudaに送って、そのまま扱えているがCの変数型はintでもfloatでもcharでも扱えるらしい。 石本さんいわくfloat32とfloat64で全然演算のパフォーマンスが違ったりするようなこともあるとのこと。最適化のときは気にしてみてみるといいかもしれない。

あとはsumのように周りのピクセル情報が必要な処理をCUDA上でうまく計算するのは少しアルゴリズム的な工夫がいるようになってくるらしい。 またループで書くと遅かったりするところをC++のテンプレートとかで展開したりというのも最適化のときにするらしい。 その辺りの特有のテクニックはまたコードを書きながら覚えていく。コードの書き方の流れは理解したのでツールチェインに関して調べる。 石本さんいわく↓の本がよかったらしい。

CUDA C プロフェッショナル プログラミング (impress top gear)

CUDA C プロフェッショナル プログラミング (impress top gear)

nvcc: NVIDIA CUDA Compiler

cuda-samplesの0_Simple/vectorAddのサンプルをbuildして、実行してみる。 zcheeさんがGithubで公開していたのでソースコード読みたい方はそちらを 0_Simple/vectorAdd

$ cd ~/cuda_samples/0_Simple/vectorAdd
$ make
/usr/local/cuda-10.0/bin/nvcc -ccbin g++ -I../../common/inc  -m64    -gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35 -gencode arch=compute_37,code=sm_37 -gencode arch=compute_50,code=sm_50 -gencode arch=compute_52,code=sm_52 -gencode arch=compute_60,code=sm_60 -gencode arch=compute_61,code=sm_61 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_75,code=sm_75 -gencode arch=compute_75,code=compute_75 -o vectorAdd.o -c vectorAdd.cu
/usr/local/cuda-10.0/bin/nvcc -ccbin g++   -m64      -gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35 -gencode arch=compute_37,code=sm_37 -gencode arch=compute_50,code=sm_50 -gencode arch=compute_52,code=sm_52 -gencode arch=compute_60,code=sm_60 -gencode arch=compute_61,code=sm_61 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_75,code=sm_75 -gencode arch=compute_75,code=compute_75 -o vectorAdd vectorAdd.o
$ make run
./vectorAdd
[Vector addition of 50000 elements]
Copy input data from the host memory to the CUDA device
CUDA kernel launch with 196 blocks of 256 threads
Copy output data from the CUDA device to the host memory
Test PASSED
Done

https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html でnvccの詳細があるので適宜みていく。 基本的にはgccのスーパーセットのようになっていて、CUDAで実行する処理はx86_64用のアセンブリではなくCUDA用の命令セットで実行ファイルが生成される。また nvlink とよばれるliinkerも含まれていてうまく呼び出しをリンクしてくれる。 このMakefileではすべてnvccでオブジェクトファイルも実行ファイルも生成しているが、CUDAを利用しているプログラム部分は nvccコンパイルしてオブジェクトファイルを作成し、呼び出し側のCのプログラムはgccとかでコンパイル、リンクすることができるらしい 詳しくはこちらを参照

ところでコードの中には __global____host__ などが含まれている。 これらは Operator と呼ばれているらしく、↓のように非常にシンプルな意味をもつ。

  • __global__: ホストから呼び出されてデバイス上で実行される関数
  • __device__: デバイスから呼び出されてデバイス上で実行される関数。ホストコードからは呼び出せない
  • __host__: ホストから呼び出されてホスト上で実行される関数(デフォルト)。__host____device__ 識別子を一緒に使うとCPU、GPUコードをともに生成。

https://http.download.nvidia.com/developer/cuda/jp/CUDA_Programming_Basics_PartII_jp.pdf

NVIDIA Visual Profiler

プロファイリングにはNVIDIA Visual Profilerを用いる。 GUIのツールだけれど、nvprofというコマンドラインがあり、基本はこれでプロファイルデータだけ作成してローカルに転送、NVIDIA Visual Profilerで閲覧したりして使える。プロファイルデータの拡張子は基本的に .nvvp とするみたい。

$ nvprof -o profile.nvvp vectorAdd

次の記事が画像つきでよくまとまっている。

vectorAddのプロファイリング結果を出してみると次のようになった。

f:id:nwpct1:20181110145310p:plain

細かい部分はまだちゃんと最適化するフェーズになってから調べていく。 まずは動くものを作ってみる。

IntelliJ CLion, cmakeの設定

CUDA Toolkitとかをcuda libraryとかも一緒に入ってくれるので、手元のmacbookでコード書くときにIDEの補完できたりしそう。 IDEに関しては自分はCLionを使っているのでcmake前提になる。 CUDAの命令セットを生成するために、CUDAの開発ではgccではなくnvcc(NVIDIA CUDA Compiler)を呼び出す必要がある。 その辺りはまだ設定してないが次の資料が役に立ちそう。

vectorAddプログラムはこれでbuildできた (macGPUはないので実行はしてない)。

cmake_minimum_required(VERSION 3.12)
project(cussim)

find_package(CUDA REQUIRED)
message("-- CUDA --")
message("Version: " ${CUDA_VERSION})
message("Library: " ${CUDA_CUDA_LIBRARY})
message("Runtime: " ${CUDA_CUDART_LIBRARY})
cuda_add_executable(vectorAdd vectorAdd.cu)

CUDA Libraryがうまく読み込めないけどとりあえずそれ以外のcmakeの機能はちゃんと使えている。 cmake build も問題なく通った。

$ cmake build
-- CUDA --
Version: 10.0
Library: CUDA_CUDA_LIBRARY-NOTFOUND
Runtime: /usr/local/cuda/lib/libcudart.dylib
-- Configuring done
-- Generating done
-- Build files have been written to: /Users/a14737/src/github.com/c-bata/cussim/build

nvidia-docker2

nvidia-docker2も触っておく。

curl -s -L https://nvidia.github.io/nvidia-docker/gpgkey | \
  sudo apt-key add -
distribution=$(. /etc/os-release;echo $ID$VERSION_ID)
curl -s -L https://nvidia.github.io/nvidia-docker/$distribution/nvidia-docker.list | \
  sudo tee /etc/apt/sources.list.d/nvidia-docker.list
sudo apt-get update
sudo apt-get install -y docker.io

libcudf(GPU DataFrame)

https://github.com/rapidsai/cudf/tree/master/libgdf もキャッチアップしておきたいので読んでおく。cudaは9.2+でrequirementsを満たしているので、g++とcmakeを確認

$ g++ --version
g++ (Ubuntu 7.3.0-27ubuntu1~18.04) 7.3.0
Copyright (C) 2017 Free Software Foundation, Inc.
This is free software; see the source for copying conditions.  There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE
$ cmake --version
cmake version 3.10.2

g++のバージョンが 5.4 よりだいぶ新しいようだがこのまま進める。 cmakeはapt-get upgradeでも3.10までしかあがらなかったので公式サイト からとってくる。

$ wget https://cmake.org/files/v3.12/cmake-3.12.4-Linux-x86_64.sh
$ sudo mkdir /opt/cmake
$ sudo ./cmake-3.12.4-Linux-x86_64.sh --skip-license --prefix=/opt/cmake
$ cat >> ~/.bash_profile <<EOF
if [ -d /opt/cmake/bin/ ]; then
    export PATH=/opt/cmake/bin/${PATH:+:${PATH}}
fi
$ source ~/.bash_profile
$ cmake --version
cmake version 3.12.4

CMake suite maintained and supported by Kitware (kitware.com/cmake).
EOF

あとc++の行列演算ライブラリ boostも使っているようなのでREADMEに従ってインストール

$ sudo apt-get install libboost-all-dev
$ cat >> ~/.bash_profile <<EOF
export ARROW_VERSION=0.10.0
export PARQUET_ARROW_VERSION=apache-arrow-$ARROW_VERSION
EOF

condaのインストール

$ conda env create --name cudf_dev --file conda/environments/dev_py35.yml
$ git clone ...
$ cd cudf
$ conda env create --name cudf_dev --file conda/environments/dev_py35.yml
$ source activate cudf_dev
$ cd cpp
$ mkdir build
$ cd build
$ cmake .. -DCMAKE_INSTALL_PREFIX=/usr/local/lib
...

[100%] Completed 'GoogleTest'
[100%] Built target GoogleTest
-- GoogleTest installed here: /home/shibata_masashi/cudf/cpp/build/googletest/install
-- Google C++ Testing Framework (Google Test) found in /home/shibata_masashi/cudf/cpp/build/googletest
-- Using Nvidia Tools Extension
-- Using legacy allocator for hash tables
-- Configuring done
CMake Error at CMakeLists.txt:163 (add_library):
  Cannot find source file:

    thirdparty/cnmem/src/cnmem.cpp

  Tried extensions .c .C .c++ .cc .cpp .cxx .cu .m .M .mm .h .hh .h++ .hm
  .hpp .hxx .in .txx


CMake Error at CMakeLists.txt:163 (add_library):
  No SOURCES given to target: rmm

-- Build files have been written to: /home/shibata_masashi/cudf/cpp/build

cnmemというNVIDIAのCUDAメモリー管理ライブラリーが必要らしい。 GitHub - NVIDIA/cnmem: A simple memory manager for CUDA designed to help Deep Learning frameworks manage memory のREADMEを見ながら入れる。

$ cd $HOME
$ cd cnmem
$ mkdir build
$ cd build
$ cmake ..
$ make
 :
-- Build files have been written to: /home/shibata_masashi/cnmem/build

header fileはあるけどライブラリーが見つからない。一旦諦める。進捗あったらまた書きます。 1回 Ubuntu16.04 でnvidia-docker2いれてコンパイルしてみたほうが早そう。

FFmpeg libavcodec

何を作るか悩ましいが今回は1つFFmpegの何らかの処理をAccelerateするようなライブラリを書いてみることをゴールにしてみる。 SSIMとかの計算をいきなりCUDAで実装するのは結構難しそうなので、素朴にCUDAで行列演算がかけそうなMSE/PSNRから始めることにする。 2つの動画を入力に受け取り、キーフレームをデコード、受け取ったAVFrameのdataを変換する。

Ubuntu 18.04 にlibav群を入れる必要があるので、 CompilationGuide/Ubuntu – FFmpeg をもとにffmpegをインストールする。 生成物をみてみるとffmpegなどのコマンドラインツールとともに、static libraryも PREFIX/lib に生成してくれている。

$ ls ../ffmpeg_build/lib/
libaom.a       libavfilter.a  libfdk-aac.a   libmp3lame.la  libpostproc.a    libvpx.a   libyasm.a
libavcodec.a   libavformat.a  libfdk-aac.la  libopus.a      libswresample.a  libx264.a  pkgconfig
libavdevice.a  libavutil.a    libmp3lame.a   libopus.la     libswscale.a     libx265.a
$ ls ../ffmpeg_build/include/
aom      libavcodec   libavformat  libswresample  libyasm-stdint.h  vpx            x265.h
fdk-aac  libavdevice  libavutil    libswscale     libyasm.h         x264.h         x265_config.h
lame     libavfilter  libpostproc  libyasm        opus              x264_config.h
$ gcc -o hello -Wall -L../ffmpeg_build/lib -I../ffmpeg_build/include -lavformat -lavcodec hello_libav.c
/tmp/ccYf66lk.o: In function `av_make_error_string':
hello_libav.c:(.text+0x24): undefined reference to `av_strerror'
/tmp/ccYf66lk.o: In function `main':
hello_libav.c:(.text+0x5e): undefined reference to `avformat_alloc_context'
hello_libav.c:(.text+0xc6): undefined reference to `avformat_open_input'
hello_libav.c:(.text+0x13d): undefined reference to `avformat_find_stream_info'
hello_libav.c:(.text+0x2c9): undefined reference to `avcodec_find_decoder'
hello_libav.c:(.text+0x3ab): undefined reference to `avcodec_alloc_context3'
hello_libav.c:(.text+0x3e6): undefined reference to `avcodec_parameters_to_context'
hello_libav.c:(.text+0x41d): undefined reference to `avcodec_open2'
hello_libav.c:(.text+0x441): undefined reference to `av_frame_alloc'
hello_libav.c:(.text+0x46e): undefined reference to `av_packet_alloc'
hello_libav.c:(.text+0x504): undefined reference to `av_packet_unref'
hello_libav.c:(.text+0x517): undefined reference to `av_read_frame'
hello_libav.c:(.text+0x53e): undefined reference to `avformat_close_input'
hello_libav.c:(.text+0x54a): undefined reference to `avformat_free_context'
hello_libav.c:(.text+0x556): undefined reference to `av_packet_free'
hello_libav.c:(.text+0x562): undefined reference to `av_frame_free'
hello_libav.c:(.text+0x56e): undefined reference to `avcodec_free_context'
/tmp/ccYf66lk.o: In function `decode_packet':
hello_libav.c:(.text+0x6de): undefined reference to `avcodec_send_packet'
hello_libav.c:(.text+0x79b): undefined reference to `avcodec_receive_frame'
hello_libav.c:(.text+0x8ae): undefined reference to `av_get_picture_type_char'
hello_libav.c:(.text+0x963): undefined reference to `av_frame_unref'
collect2: error: ld returned 1 exit status

見つからない関数群があるみたいなので、libavcodeclibavformatのバージョンが違うのかなと思いつつ、一応symbolを見てみる

$ grep -nr "avformat_alloc_context" ./include/
./include/libavformat/avformat.h:55: * AVFormatContext, use avformat_alloc_context() (some functions, like
./include/libavformat/avformat.h:122: * avformat_alloc_context() and do some tweaking on it before passing it to
./include/libavformat/avformat.h:198: * avformat_alloc_context() to create a muxing context. The caller then sets up
./include/libavformat/avformat.h:1329: * avformat_alloc_context() to create an AVFormatContext.
./include/libavformat/avformat.h:1339:     * A class for logging and @ref avoptions. Set by avformat_alloc_context().
./include/libavformat/avformat.h:2118:AVFormatContext *avformat_alloc_context(void);
 :
$ nm ~/ffmpeg_build/lib/libavformat.a | grep avformat_alloc_context
                 U avformat_alloc_context
                 U avformat_alloc_context
                 :

ありそう... リンクの順番ミスってる?

$ gcc -c -o hello.o -Wall -I../ffmpeg_build/include hello_libav.c

ちょっとわからないので一旦ここまでで放置。 進展あればまた追記します。

おまけ: NVIDIA Tesla K80

今回はNVIDIA Tesla K80を利用した。 K80に関する自分用のメモを残しておく。

  • CUDA Core: 4992(2496x2) 562MHz
  • Memory 24GiB. 帯域幅 480GB/s (240x2). PCI Express Gen3
    • この数値はGPU上のメモリーの読み書きの速度。CPUのメモリからの読み出しが30GB/sとかだったりするから、Host to Deviceの転送スピードとしてはこっちだけ早くても仕方ないらしい。
    • 石本さんいわくGPU to GPUが早かったりとかはするかも。前に石元さんが試したときは普通に転送すると9GB/sとか10GB/s。Cでmemcpyすると9-10GB/secondsぐらいだからおそらくCPU側のメモリの読み以下期速度がボトルネックだろうとのこと。
  • GCP台湾リージョンで使える唯一のGPU
  • NVENCは H.264 (AVCHD) YUV 4:2:0 にのみ対応している

その他のGPUについては https://ja.wikipedia.org/wiki/NVIDIA_Tesla などによくまとまっていた。

RAW Socket / BPF(Berkeley Packet Filter)を用いたパケットキャプチャーツールの実装

パケットキャプチャーツールは、ネットワークを流れるすべてのパケットを受け取り解析します。 NIC(Network Interface Card)のほとんどはプロミスキャスモードとよばれるモードをサポートしており、これを有効にすることでアドレスにかかわらずNICはすべてのパケットをホストに渡します。 ソフトウェアとハードウェアが連携して動作するため、扱っているレイヤーが低く環境によってInterfaceに差異があります。

tcpdumpの開発者によってつくられたlibpcapというライブラリはUNIXのシステムの差異を吸収します。またWindowsにもWinPcapという名前で移植されています。 もしパケットキャプチャーを作る際にはlibpcapを利用することが一般的かと思いますが、今回は勉強も兼ねて LinuxmacOS で動作するパケットキャプチャーをlibpcapを使わずに1からC言語で実装してみました。

※ BPF VM(Berkeley Packet Filter Virtual Machine)によるFilteringの仕組みには今回は触れません。

f:id:nwpct1:20181028135417g:plain

目次

xpcap のソースコード(Github)

github.com

最近作りたいなと思っているパケットキャプチャー関連のソフトわがありそちらはGo言語で実装しているのですが、せっかくなら移植性を考えてPure Goで実装したいと思っています。こういったレイヤーのプログラムをいきなりCのAPIがベースにあってそれをGoで書くとドキュメントを追うのも大変なので、C言語でまずは書いてみたものがこちらです。

プロトコルは今のところARPIPv4IPv6TCPUDP、ICMPに対応していて、Ethernet Frameからパースした結果を標準出力に書き出します。

RAW SOCKETを用いたキャプチャー (Linux)

LinuxMACアドレスEthernet Frameのヘッダー情報までプログラムで扱うには、RAW Socketが必要です。 ソケットディスクリプタを取得する際には、アドレスファミリーとして AF_PACKET 、ソケットタイプとして SOCK_RAW そして第3引数のprotocolには htons(EATH_P_ALL) を指定します。全部を説明すると長くなるので手順と呼び出さないといけない関数を次に示します。 xpcapのソースコードと合わせてご覧ください。

  1. socket() ディスクリプタの取得
    • int soc = socket(AF_PACKET, SOCK_RAW, htons(ETH_P_ALL)))
  2. en0 などのインターフェイス名を指定してインターフェイスの情報を取得
    • ioctl(soc, SIOCGIFINDEX, &if_req)
  3. ソケットディスクリプターをインターフェイスにバインド
    • bind(soc, (struct sockaddr *) &sa, sizeof(sa))
  4. インターフェイスのフラグを取得
    • ioctl(soc, SIOCGIFFLAGS, &if_req)
  5. プロミスキャスモードを有効にし、インターフェイスをUP(動作中)にする
    • ioctl(soc, SIOCSIFFLAGS, &if_req)

これで準備が完了です。あとは selectepoll でソケットディスクリプターへの書き込みを監視しready担った状態で recv(2) で読み出せばOKです。

struct timeval timeout;
fd_set mask;
int width, len, ready;
while (g_gotsig == 0) {
    FD_ZERO(&mask);
    FD_SET(soc, &mask);
    width = doc + 1;

    timeout.tv_sec = 8;
    timeout.tv_usec = 0;
    ready = select(width, &mask, NULL, NULL, &timeout);
    if (ready == -1) {
        perror("select");
        break;
    } else if (ready == 0) {
        fprintf(stderr, "select timeout");
        break;
    }

    if (FD_ISSET(sniffer->fd, &mask)){
        if ((len = recv(soc, buffer, >buf_len, 0)) == -1){
            perror("recv:");
            return -1;
        }
    }
}

自分は Linuxネットワークプログラミングバイブル で勉強しましたが、この書籍以外にもLinuxで動くRAW SOCKETを使ったシンプルなパケットキャプチャーの作り方を解説している資料は多くあります。一方でBSD系のOSではアドレスファミリーとして AF_PACKET を指定できません。BSD系のOSでEthernet frameを読み出す方法を確認しましょう。

BPF(Berkeley Packet Filter)によるキャプチャー (macOS, BSD系)

これらはBPF(Berkeley Packet Filter)という仕組みを使う必要があります。 BPFにはBPF Virtual Machineという仕組みを使ってパケットをKernel側でフィルタリングすることで必要ないものまでユーザー空間に移さずオーバーヘッドを減らす仕組みのようです。読み出しには BPFデバイスというのを用います。ひとまずすべてキャプチャーするならBPF VMについては気にする必要はありません。

BPFデバイスは、 /dev/bpf* に存在します。これらを順にopenしながら、使用可能なBPFデバイスを探さなくてはいけません。

$ ls /dev/bpf?
/dev/bpf0 /dev/bpf1 /dev/bpf2 /dev/bpf3 /dev/bpf4 /dev/bpf5 /dev/bpf6 /dev/bpf7 /dev/bpf8 /dev/bpf9

手元では bpf255 ぐらいまで存在しますが、google/gopacketなどの実装では99までチェックしているようです。 NICの数以上に必要になるケースはほとんどなさそうなので99個は十分に余裕を持った値なんだと思います。

gopacket/bsd_bpf_sniffer.go at a35e09f9f224786863ce609de910bc82fc4d4faf · google/gopacket · GitHub

BPFデバイスが決まったらOpenします。その後次のような手順が準備に必要になります。

  1. bpfデバイスのopen
    • fd = open(params.device, O_RDWR)
  2. バッファ長の設定 or 取得。BIOCSBLEN の変更は、BPFデバイスNICアサインする BIOCSETIF より先に呼び出される必要があるので注意してください。これになかなか気づかず結構はまってしまいました。
    • ioctl(fd, BIOCSBLEN, &params.buf_len) : 設定
    • ioctl(fd, BIOCGBLEN, &params.buf_len) : 取得
  3. BPFデバイスとネットワークインターフェイスをバインド
    • ioctl(fd, BIOCSETIF, &if_req)
  4. プロミスキャスモードの有効化
    • ioctl(fd, BIOCPROMISC, NULL)

こちらはデバイスファイルなので recv(2) ではなく read(2) で読み出します。 読み出すとイーサネットのフレームではなくBPFパケットというものにくるまれています。 ヘッダーをパースするとデータ長が乗っているため、それをもとに次のBPFパケットの位置を求めてパースを繰り返していきます。

typedef struct {
    int fd;
    char device[11];
    unsigned int buf_len;
    char *buffer;
    unsigned int last_read_len;
    unsigned int read_bytes_consumed;
} Sniffer;

int
parse_bpf_packets(Sniffer *sniffer, CapturedInfo *info)
{
    if (sniffer->read_bytes_consumed + sizeof(sniffer->buffer) >= sniffer->last_read_len) {
        return 0;
    }

    info->bpf_hdr = (struct bpf_hdr*)((long)sniffer->buffer + (long)sniffer->read_bytes_consumed);
    info->data = sniffer->buffer + (long)sniffer->read_bytes_consumed + info->bpf_hdr->bh_hdrlen;
    sniffer->read_bytes_consumed += BPF_WORDALIGN(info->bpf_hdr->bh_hdrlen + info->bpf_hdr->bh_caplen);
    return info->bpf_hdr->bh_datalen;
}

あとはごりごりパースしていくのですが、そこはプラットフォームに変わらず同じです。 ゴリゴリ実装していくだけで解説してもしかたないのでマスタリングTCP/IPなどを頼りにソースコードを読んでみてください。

実行方法

build.sh でビルドできます。Vagrantfileも用意しているのでLinuxで試したいmacOSユーザーの方はご利用ください。実行結果は次のような感じです。

$ ./build.sh
$ ./xpcap en0 -v
device = en0, verbose = 1, port = 0

================================================================================
[TCP6]
ether_header--------------------------------------------------------------------
ether_dhost = XX:XX:XX:XX:XX:XX
ether_shost = XX:XX:XX:XX:XX:XX
ether_type = 86DD(IPv6)
ip6-----------------------------------------------------------------------------
ip6_vfc = 96
ip6_flow = 2363892320
ip6_plen = 15104
(TCP), ip6_hlim = 56
ip6_src = xxxx:xxxx:xxxx:x::xxxx:xxxx
ip6_dst = yyyy:yy:yyyy:yyyy:yyyy:yyyy:yyyy:yyyy
tcphdr--------------------------------------------------------------------------
source: 47873
destination: 59083
sequence number: 1148644729
ack number = 2897299570
data offset = 5, control flag = 24, window = 49152, checksum = 54057, urgent pointer = 0
data----------------------------------------------------------------------------
00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00    ..something..data..
================================================================================

================================================================================
[ARP]
ether_header--------------------------------------------------------------------
ether_dhost = XX:XX:XX:XX:XX:XX
ether_shost = XX:XX:XX:XX:XX:XX
ether_type = 806(Address resolution)
ether_arp-----------------------------------------------------------------------
arp_hrd = 1(Ethernet 10/100Mbps.), arp_pro = 2048(IP)
arp_hln = 6, arp_pln = 4, arp_op = 1(ARP request.)
arp_sha = 34:76:C5:77:5D:4C
arp_spa = 192.168.0.1
arp_tha = 00:00:00:00:00:00
arp_tpa = 192.168.0.8
================================================================================

================================================================================
[UDP]
ether_header--------------------------------------------------------------------
ether_dhost = XX:XX:XX:XX:XX:XX
ether_shost = XX:XX:XX:XX:XX:XX
ether_type = 800(IP)
ip------------------------------------------------------------------------------
ip_v = 4, ip_hl = 5, ip_tos = 0, ip_len = 149
ip_id = 29282, ip_off = 0, 0
ip_ttl = 255, ip_p = 17(UDP), ip_sum = 42831
ip_src = yyy.yyy.yyy.yyy
ip_dst = xxx.xxx.xxx.xxx
udphdr--------------------------------------------------------------------------
source = 5353, dest = 5353
len = 129, check = 38825
data----------------------------------------------------------------------------
00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00    ..something..data..
================================================================================

参考ソースコード

困ったときは次のコードが参考になりました。

またLinuxネットワークプログラミングバイブルはかなりおすすめの書籍です。

Linuxネットワークプログラミングバイブル

Linuxネットワークプログラミングバイブル

実践 パケット解析 第3版 ―Wiresharkを使ったトラブルシューティング

実践 パケット解析 第3版 ―Wiresharkを使ったトラブルシューティング