c-bata web

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

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を使ったトラブルシューティング

Google Translate APIを使ったSphinxドキュメントの自動翻訳

多言語への翻訳は大変な作業ですが、近年は機械翻訳の精度も上がってきました。 ふと思いついて .po 形式の翻訳ファイルをGoogle翻訳を通して自動で入力するスクリプト を作ったのですが、サクッと書いた割に予想以上に便利で料金も思ったより安かったので記事にしました。また実際に自分が公開している日本語で書かれたSphinxの資料を、このスクリプトを使って英語に翻訳してみます。

f:id:nwpct1:20181029112031p:plain

追記: ライセンスについて

id:beatdjam さんのコメントが気になったので共有です。 自分もGoogle Translate APIのドキュメントを読み返してみますが、利用される方も確認してからご利用ください。

以前こういった事例もあったので、OSSで利用することを推奨して良いのか心配。ドキュメントだけなら平気なのかな https://anond.hatelabo.jp/20170225195916

作ったもの

.po 形式の翻訳ファイルをパースし好きな言語にGoogle Translate APIを用いて翻訳するスクリプトを用意しました。 実行には google-cloud-translate とGCP service accountが必要です。

$ pip install --upgrade google-cloud-translate
$ export GOOGLE_SERVICE_ACCOUNT_JSON=/path/to/service-account-credential.json
$ python translate_po.py --help
usage: translate_po.py [-h] [--lang LANG] [--currency CURRENCY] filepath

positional arguments:
  filepath

optional arguments:
  -h, --help           show this help message and exit
  --lang LANG          target language (default: "ja")
  --currency CURRENCY  dollar per your currency. (default currency is yen: 111.90)

現状はとりあえずファイルの上書きオプションなどは用意せず、stdoutに書き出すようにしています。 Google Translate APIは、100万文字あたり20ドルかかります。 本一冊とかになると数百円かそれ以上かかりそうですが、手元の文章を翻訳したいなどの用途なら数十円に収まることがほとんどです。 ちなみにマルチバイト文字でも1文字は1文字としてカウントしてくれるようなので、日本語から英語の翻訳などは比較的お得です。 Google Translate APIに投げたテキストの文字数からかかった金額も算出し表示するようにしています。

$ python translate_po.py ./po/index.po 1>./po/index_ja.po
Cost: 2.1417659999999996 yen

また翻訳結果はキャッシュしていて、実行したディレクトリの直下に json ファイルを書き出します。 なので2回目の実行は、キャッシュが効きお金を節約できます。

$ python translate_po.py ./po/index.po 1>./po/index_ja.po
Cost: 0 yen

ソースコードGithubで公開しています。

github.com

実際に翻訳してみる

Webアプリケーションフレームワークの作り方 in Python — c-bata.link (Githubはこちら) はSphinxで書かれた日本語の資料です。 今回はこちらを英語に翻訳していきます。Sphinxのドキュメントの国際化の方法は次のページに非常によくまとまっています。

まず sphinx-intl をインストールします。

$ pip install sphinx-intl
$ vim source/conf.py
# add following settings
# locale_dirs = ['locale/']
# gettext_compact = False
$ make gettext
$ ls build/locale/
index.pot      kobin.pot      middleware.pot request.pot    response.pot   routing.pot    server.pot     sphinx.pot     template.pot   wsgi.pot

potファイルができました。今回は日本語から英語に翻訳するので、次のようにします。

$ sphinx-intl update -p build/locale -l ja
Create: source/locale/ja/LC_MESSAGES/kobin.po
Create: source/locale/ja/LC_MESSAGES/template.po
Create: source/locale/ja/LC_MESSAGES/middleware.po
Create: source/locale/ja/LC_MESSAGES/sphinx.po
Create: source/locale/ja/LC_MESSAGES/request.po
Create: source/locale/ja/LC_MESSAGES/routing.po
Create: source/locale/ja/LC_MESSAGES/wsgi.po
Create: source/locale/ja/LC_MESSAGES/response.po
Create: source/locale/ja/LC_MESSAGES/index.po
Create: source/locale/ja/LC_MESSAGES/server.po

poファイルが出来上がったら変換をかけます。 このスクリプトは今のところ上書き用のオプションを用意していないので、一度stdoutをファイルに書き出して置き換える必要があります。 いくつかファイルがあるので変換用のスクリプトを用意しました。

$ cat > translate.sh <<EOF
#!/bin/bash
 function translate {
  for f in ./ja/LC_MESSAGES/*.po; do
    python translate_po.py --lang en $f 1>${f%.po}_en.po
    mv ${f%.po}_en.po $f;
  done; 
}
 translate
EOF
$ chmod +x ./translate.sh
$ ./translate.sh
Cost: 3.6904620000000006 yen
Cost: 0.024617999999999998 yen
Cost: 1.60017 yen
Cost: 4.728894 yen
Cost: 3.8784539999999996 yen
Cost: 5.8188 yen
Cost: 1.087668 yen
Cost: 1.4009880000000001 yen
Cost: 7.00494 yen

トータル30円くらいかかりました。翻訳精度を考えるとすごくお得に感じます。 最後はこれをbuildしてみましょう。

$ make -e SPHINXOPTS="-D language='ja'" html
$ open build/html/index.html

結果は次のような感じです。

f:id:nwpct1:20181028231948p:plain

f:id:nwpct1:20181028232035p:plain

f:id:nwpct1:20181028232044p:plain

reSTのリンクが壊れたり、いくつか変な文字が混ざっていたりはしますが予想以上にそれっぽくなりました。

エキスパートPythonプログラミング 改訂2版 (アスキードワンゴ)

エキスパートPythonプログラミング 改訂2版 (アスキードワンゴ)