AI」カテゴリーアーカイブ

Android UserLAnd で PyTorch を使う。C++ API

Android スマートフォンの上に PyTorch の開発環境を作ってみました。かなり遅いですが、電車の中で書いたコードのビルドテストをしたり、簡単な実験をするくらいには使えそうです。

以下作業メモです。

使用した端末は Essential Phone PH-1 (Snapdragon 835, RAM 4GB, Android 10)。もちろん C++ だけでなく Python からも使えます。

●UserLAnd の準備

UserLAnd を使うと Android 上に Linux 環境を作ることができます。root 不要でただの Android アプリとして動きます。今回は最小構成にするため Ubuntu + SSH Terminal を選択しました。

●PyTorch のビルド

UserLAnd の Ubuntu を起動し Terminal で作業します。または PC 等から ssh で接続することもできます。

先に更新しておきます。

$ sudo apt update
$ sudo apt upgrade -y

必要なソフトのインストール

$ sudo apt install  git cmake g++
$ sudo apt install  libopenblas-dev

Python3 の準備

$ sudo apt install  python3-dev python3-setuptools
$ sudo apt install  python3-numpy python3-cffi
$ sudo apt install  python3-yaml python3-wheel

pytorch の準備。展開フォルダは仮に ~/pytorch とします。

$ cd
$ git clone --recursive https://github.com/pytorch/pytorch
$ cd pytorch
$ git submodule sync
$ git submodule update --init --recursive

ビルドします。複数スレッド使うとメモリ不足で落ちるので MAX_JOBS=1 を指定します。

$ cd ~/pytorch
$ USE_CUDA=0 USE_CUDNN=0 BUILD_TEST=0 USE_MKLDNN=0 USE_DISTRIBUTED=0 MAX_JOBS=1 python3 setup.py build_clib

Python 向けにビルドする場合は最後の “build_clib” を “install” にしてください。

かなり時間がかかるので電源に接続したまましばらく待ちます。

もし途中で↓のようなエラーが出た場合はコードを一部修正します。

... 8x8-dq-aarch64-neon.S:662: Error: operand mismatch -- `mov V8.4s,V9.4s'

“aten/src/ATen/native/quantized/cpu/qnnpack/src/q8gemm/8×8-dq-aarch64-neon.S” の 662行~669行あたりにある

    MOV V8.4s, V9.4s
    MOV v10.4s, v11.4s
    MOV v12.4s, V13.4s
    MOV V14.4s, V15.4s
    MOV V16.4s, V17.4s
    MOV V18.4s, V19.4s
    MOV V20.4s, V21.4s
    MOV V22.4s, V23.4s

これらの行を下のように変更します。

    MOV V8.16b, V9.16b
    MOV v10.16b, v11.16b
    MOV v12.16b, V13.16b
    MOV V14.16b, V15.16b
    MOV V16.16b, V17.16b
    MOV V18.16b, V19.16b
    MOV V20.16b, V21.16b
    MOV V22.16b, V23.16b

AArch64 の NEON にはもともとレジスタ同士の move 命令がありません。mov vdest, vsrc はアセンブラ内部で「 or vdest, vsrc, vsrc 」に置き換えられます。浮動小数点数で論理演算はできないので、アセンブラによっては 4s (単精度 32bit x4 = 128bit) の型指定がエラーになってしまうのだと考えられます。これは 16b (byte 8bit x16 = 128bit) に変更すれば OK です。

● C++ から利用する

前回も書いたように ~/pytorch/torch 以下が libtorch 相当になります。

include path に

~/pytorch/torch/include
~/pytorch/torch/include/torch/csrc/api/include

link path に

~/pytorch/torch/lib

を指定してコンパイルします。必要なライブラリは libc10 と libtorch_cpu です。コンパイルオプションの指定例は下記の通り。

$ clang -L~/pytorch/torch/lib -lc10 -ltorch_cpu  torchtest.cpp -o torchtest

実行時は LD_LIBRARY_PATH が必要です。

$ export LD_LIBRARY_PATH=~/pytorch/torch/lib

一度ビルドしておけば他の AArch64 Android の UserLAnd 上でも使えます。Raspberry Pi 4 の Ubuntu Server (AArch64) でも動きました。

● Python で使う

Python の場合は “build_clib” の代わりに “install” を使います。途中でエラーが出ますが問題ありません。(“bdist_wheel” は途中で止まるので保留中)

$ USE_CUDA=0 USE_CUDNN=0 BUILD_TEST=0 USE_MKLDNN=0 USE_DISTRIBUTED=0 MAX_JOBS=1 python3 setup.py install

これで import torch できるようになります。torchvision が必要な場合も同じようにビルドしておきます。

$ sudo apt install zlib1g-dev libjpeg-dev
$ git clone https://github.com/pytorch/vision.git
$ cd vision
$ MAX_JOBS=1 python3 setup.py install

関連ページ
HYPERでんち: Deep Learning

関連エントリ
RADEON (ROCm) で PyTorch を使う。C++ API
Jetson Nano で TensorFlow の C 言語 API を使う
Android UserLAnd の更新と VNC 画面設定
UserLAnd : Android 9.0 で Ctrl + SPACE を使えるようにする
Android Termux で日本語入力を行う / UserLAnd との併用
Android 9.0 と Bluetooth Keyboard による日本語入力
Android で動く Linux 環境 UserLAnd が XServer XSDL に対応
Oculus Go を文章書き&開発マシンにする
UserLAnd とブラウザ
Android 上の開発環境と UserLAnd
OS の中の Linux (WSL/Chrome OS/Android UserLAnd)
ARM CPU 上の開発環境とコンパイル時間の比較 (2) Pixel 3/UserLAnd

RADEON (ROCm) で PyTorch を使う。C++ API

Deep Learning フレームワークのほとんどが Python を使っています。C++ など他の言語も使えますが、Python と同じ使い方ができるものはあまり多くありません。

例えば CNTK の場合 Keras の backend として使うとこんな感じです。

# Keras
model= models.Sequential([
    layers.Dense( 64, input_shape=(784,), activation='relu' ),
    ~
  ])

CNTK の Python API を使った場合↓

# CNTK
xinput= cntk.input_variable( 784, np.float32 )
model= layers.Sequential([
    layers.Dense( 64, activation=cntk.relu ),
    ~
  ])

CNTK の C++ の場合 Parameter を確保して複数の Function に分解する必要があるのでこうなります。↓

// CNTK + C++
auto device= CNTK::DeviceDescriptor::UseDefaultDevice();
auto xinput= CNTK::InputVariable( { 784 }, CNTK::DataType::Float, L"xinput" );
auto weight= CNTK::Parameter( { 64, 784 }, CNTK::DataType::Float, CNTK::HeNormaliInitializer(), device );
auto bias= CNTK::Parameter( { 64 }, CNTK::DataType::Float, 0.0f, device );
auto x= CNTK::ReLU( CNTK::Plus( bias, CNTK::Times( weight, xinput ) ) );
~

PyTorch の場合は C++ でも Python と同じ書き方で同じ API なので覚えることが少なくて済みます。

// PyTorch C++
class ModelImpl : public torch::nn::Module {
    torch::nn::Linear  fc0;
public:
    ModelImpl()
    {
        fc0= register_module( "fc0", torch::nn::Linear( 784, 64 ) );
	~
    }
    torch::Tensor  forward( const torch::Tensor& x )
    {
        ~
        return  torch::relu( fc0( x ) );
    }
};
TORCH_MODULE(Model);

AMD RADEON を使用する場合の選択肢はいくつかあります。Python の場合は Keras + PlaidML が最も簡単で Windows/macOS でも使用できます。Linux の場合は ROCm が使えるため PlaidML 以外の選択肢が増えます。今回は C++ API を使いたいので Linux 上で ROCm を使用しました。

以下 RADEON で PyTorch (C++ API) を使うための作業メモです。

● ROCm の install

使用した環境は RADEON RX Vega 64 (gfx900) + Ubuntu 18.04 LTS です。

ROCm Documentation: Installation Guide

上のページを見ると Supported OS に 18.04.3 (Kernel 5.0) と書かれているため最初に Kernel を更新します。下記ページを参考にさせていただきました。

virtualiment: Ubuntu18.04にカーネル 5.0 をインストールする手順

カーネルの更新

$ sudo apt install --install-recommends linux-generic-hwe-18.04 xserver-xorg-hwe-18.04 
$ sudo reboot

次に手順通りに ROCm を入れます。こちらは下記のページを参考にさせていただきました。

ニートが始めるUE4開発日誌: RX470で機械学習ことはじめ その3 ~ROCmとTensorFlowのインストール~

install 前の更新とライブラリインストール

$ sudo apt update
$ sudo apt dist-upgrade
$ sudo apt install libnuma-dev
$ sudo reboot

リポジトリ追加と ROCm のインストール

$ wget http://repo.radeon.com/rocm/apt/debian/rocm.gpg.key
$ sudo apt-key add rocm.gpg.key
$ echo 'deb [arch=amd64] http://repo.radeon.com/rocm/apt/debian/ xenial main' | \
   sudo tee /etc/apt/sources.list.d/rocm.list
$ sudo apt install rocm-dkms
$ sudo usermod -a -G video $LOGNAME
$ sudo reboot

● PyTorch のビルド

PyTorch を ROCm でビルドします。ビルド手順は下記ページを参考にさせていただきました。

Lernapparat: Building PyTorch on ROCm

事前の準備

$ sudo apt install git cmake
$ sudo apt install python3-pip
$ pip3 install setuptools numpy

ビルドに必要な ROCm のパッケージを入れます。

$ sudo apt install rocm-libs miopen-hip rccl roctracer-dev

cmake のパッチのため下記のスクリプトを使います。

#!/bin/bash
for fn in $(find /opt/rocm/ -name \*.cmake ); do
  sudo sed --in-place='~' 's/find_dependency(hip)/find_dependency(HIP)/' $fn
done

上記の内容をファイル(replace_script.sh)に保存してから下記のように実行します。

$ bash ./replace_script.sh

PyTorch を checkout します。仮に展開場所を ~/pytorch とします。

$ git clone --recursive http://github.com/pytorch/pytorch

pytorch フォルダで下記のコマンドを実行します。

$ cd ~/pytorch
$ python3 tools/amd_build/build_amd.py

ビルドします。

$ cd ~/pytorch
$ RCCL_DIR=/opt/rocm/rccl/lib/cmake/rccl/ PYTORCH_ROCM_ARCH=gfx900 hip_DIR=/opt/rocm/hip/cmake/ USE_NVCC=OFF BUILD_CAFFE2_OPS=0 PATH=/usr/lib/ccache/:$PATH USE_CUDA=OFF python3 setup.py bdist_wheel

しばらくしたら Python 向けのパッケージが ~/pytorch/dist/*.whl にできます。C++ 向けの libtorch 相当が ~/pytorch/torch 以下になります。

include path は下記の通り。

~/pytorch/torch/include
~/pytorch/torch/include/torch/csrc/api/include

library は下記の場所に入るので、ビルド時の libpath と実行時の LD_LIBRARY_PATH に追加します。

~/pytorch/torch/lib

Link 時のライブラリ指定はこんな感じ。

-lc10 -ltorch -lc10_hip -ltorch_hip -ltorch_cpu -lc10

これで CUDA 向けのコードがそのまま RADEON で動くようになりました。ROCm (hip) でも PyTorch の API は cuda のままです。C++ API でも to( torch::kCUDA ) で OK。

import torch
for di in range(torch.cuda.device_count()):
    print( di, torch.cuda.get_device_name(di), torch.cuda.get_device_capability(di) )

例えば Python で↑のコードを実行すると次のように表示されます。

0 Vega 10 XT [Radeon RX Vega 64] (3, 0)

起動時に少々待たされますが起動してしまえば十分速いです。BatchSize は 64 以上がおすすめです。

関連ページ
HYPERでんち: Deep Learning

関連エントリ
Jetson Nano で TensorFlow の C 言語 API を使う

Jetson Nano で TensorFlow の C 言語 API を使う

メモです。参考にしたページは下記の通り。TensorFlow の C言語用ライブラリを作るために実機上でビルドしています。

TensorFlow: Build from source
Building Tensorflow 1.13 on Jetson Xavier
JK Jung’s blog: Building TensorFlow 1.12.2 on Jetson Nano

bazel の install ができれば、あとは PC と同じように source からビルドできます。bazel の build も上記ページのスクリプトをそのまま利用させていただきました。

(1) 予め Jetson Nano に swap を設定しておきます。

Jetson Nano関係のTIPSまとめ
Setting up Jetson Nano: The Basics

(2) bazel を install します。

bazel の version は このページ下の表 に従い TensorFlow に合わせる必要があります。TensorFlow r1.12 の場合は 0.15.2 を使用します。こちら の script をそのまま利用させていただきました。

$ git clone https://github.com/jkjung-avt/jetson_nano.git
$ cd jetson_nano
$ ./install_bazel-0.15.2.sh

(3) TensorFlow の build

TensorFlow のソースを build します。ここ の情報に従い、若干ファイルを修正する必要があるようです。こちら の script を使うとビルド時にパッチを当ててくれます。

今回は C言語用 lib が欲しいので install_tensorflow-1.12.2.sh を下記のように修正します。

・「sudo apt-get」「sudo pip3」「bazel-bin」で始まる行をすべてコメントアウト。
・「bazel build」の実行を下記のように修正

bazel build --config=opt \
	    --config=cuda \
	    --local_resources=2048.0,1.0,1.0 \
            //tensorflow/tools/pip_package:build_pip_package

bazel build --config=opt \
            --config=cuda \
	    --config=monolithic \
	    --local_resources=2048.0,1.0,1.0 \
	    //tensorflow:libtensorflow.so

あとは script を実行すると $HOME/src/tensorflow-1.12.2 以下でビルドが行われます。

$ ./install_tensorflow-1.12.2.sh

途中でエラーが出ても、$HOME/src/tensorflow-1.12.2/bazel-bin/tensorflow 以下に libtensorflow.so が出来ていれば成功です。lib とヘッダファイルは下記の場所にあるので、必要に応じて別の場所にコピーします。

lib:     $HOME/src/tensorflow-1.12.2/bazel-bin/tensorflow/libtensorflow.so
include: $HOME/src/tensorflow-1.12.2/tensorflow/c

あとは下記のような感じで使えるはずです。

#include 
~

auto* graph= TF_NewGraph();
auto* option= TF_NewSessionOptions();
auto* status= TF_NewStatus();
auto* session= TF_NewSession( graph, option, status );
~

関連ページ
HYPERでんち: NVIDIA Jetson Nano

関連エントリ
Jetson Nano / Clang の Version とコンパイル速度の比較

Snapdragon 845 ARMv8.2A 半精度 fp16 演算命令を使ってみる / Deep Learning 命令

Snapdragon 845 の Kryo 385 (Cortex-A75) は ARMv8.2 の拡張命令である半精度浮動小数点演算に対応しています。半精度浮動小数点数は HDR テクスチャなど GPU ではおなじみで、符号 1bit 指数部 5bit 仮数部 10bit の合計 16bit で表現されます。

CPU でもこれまで単精度と半精度 (fp32 と fp16) の相互変換が可能でした。X86/X64 では F16C 命令 (vcvtph2ps/vcvtps2ph) がありますし、ARM では以前試したように Cortex-A9 以降で変換命令が追加されています。変換だけなのでメモリアクセスは速くなるものの演算速度は特に変わりません。

ARMv8.2 ではオプションの拡張命令 FPHP, SIMDHP が新設され、対応していれば 16bit 半精度のまま演算ができるようになりました。128bit の SIMD(NEON) なら同時に 8個の積和演算を行うので、ピークの演算速度は単精度の倍になる計算です。

fmla  v0.2d, v1.2d, v2.2d   ; 倍精度 64bit x2
fmla  v0.4s, v1.4s, v2.4s   ; 単精度 32bit x4
fmla  v0.8h, v1.8h, v2.8h   ; 半精度 16bit x8

新しい vfpbench で対応したので実際に計測してみたのがこちらです。8 core で 8 thread 並列時の値です。

half fp16 single fp32 double fp64
Snapdragon 845 ARMv8.2A 277.7 GFLOPS 138.4 GFLOPS 68.7 GFLOPS
Snapdragon 835 ARMv8.0A — GFLOPS 129.5 GFLOPS 67.3 GFLOPS

・GFLOPS の値が大きい方が高速

予想通りほぼ fp32 の倍の値になっています。なお big/little core を個別計測した結果の合計なので、全 core 同時に走らせた場合はもう少し数値は下がるものと思われます。big, little それぞれの値を表にすると下記の通り。

half fp16 singlel fp32 double fp64
little big little big little big
Snapdragon 845 (1.77GHz + 2.80GHz) 108.9 168.8 54.0 84.4 27.3 41.5
Snapdragon 835 (1.90GHz + 2.45GHz) 59.3 70.2 29.6 37.7

Cortex-A75 の FP/SIMD pipe は 2本ありますが、命令単位で調べると 64bit (4h) 時は IPC=2、128bit (8h) 時は IPC=1 なのでそれぞれの pipe は 64bit であることがわかります。

Deep Learning 用の命令としては他にも 8bit 積和演算である Dot Product (dotprod) 拡張命令があります。これも ARMv8.2A のオプションで、下記ような整数 Int8 の 4乗算と 4加算を 4並列で行うことができます。AVX512VNNI や GeForce RTX (Turing) などの Int8 命令によく似ています。

udot  v0.4s, v1.16b, v2.16b
sdot  v0.4s, v1.16b, v2.16b

32bit += 8bit * 8bit + 8bit * 8bit + 8bit * 8bit + 8bit * 8bit
32bit += 8bit * 8bit + 8bit * 8bit + 8bit * 8bit + 8bit * 8bit
32bit += 8bit * 8bit + 8bit * 8bit + 8bit * 8bit + 8bit * 8bit
32bit += 8bit * 8bit + 8bit * 8bit + 8bit * 8bit + 8bit * 8bit

fp16 の倍なので計算上は 500 GOPS を超えるのですが、残念ながら Snapdragon 845 の Kryo 385 では対応していませんでした。

詳細なログはこちら
Pixel 3 Snapdragon 845 Kryo 385 2.8GHz x4 + 1.77GHz x4 ARM64 (AArch64) Android 9.0

vfpbench のログ(一部)

ARCH: ARMv8.2A
FPU : ASIMD(AArch64 NEON) FPHP ASIMDHP
Name: Qualcomm Technologies, Inc SDM845

CPU Thread:  8
CPU Core  :  8
CPU Group :  2
  Group 0: Thread= 4  Clock=1.766400 GHz  (mask:f)
  Group 1: Thread= 4  Clock=2.803200 GHz  (mask:f0)
NEON  : yes
FMA   : yes
FPHP  : yes
SIMDHP: yes

Total:
SingleThread HP max:   71.675 GFLOPS
SingleThread SP max:   35.892 GFLOPS
SingleThread DP max:   17.940 GFLOPS
MultiThread  HP max:  277.711 GFLOPS
MultiThread  SP max:  138.445 GFLOPS
MultiThread  DP max:   68.745 GFLOPS

Group 0:  Thread=4  Clock=1.766400 GHz  (mask:f)
  SingleThread HP max:   27.426 GFLOPS
  SingleThread SP max:   13.683 GFLOPS
  SingleThread DP max:    6.851 GFLOPS
  MultiThread  HP max:  108.928 GFLOPS
  MultiThread  SP max:   54.046 GFLOPS
  MultiThread  DP max:   27.273 GFLOPS

Group 1:  Thread=4  Clock=2.803200 GHz  (mask:f0)
  SingleThread HP max:   44.248 GFLOPS
  SingleThread SP max:   22.209 GFLOPS
  SingleThread DP max:   11.090 GFLOPS
  MultiThread  HP max:  168.783 GFLOPS
  MultiThread  SP max:   84.400 GFLOPS
  MultiThread  DP max:   41.472 GFLOPS

関連ページ
VFP Benchmark Log 計測結果まとめ

関連エントリ
Snapdragon 835 と 845 のコンパイル時間の比較&浮動小数点演算能力
Snapdragon 845 の浮動小数点演算速度
ARM CPU の浮動小数点演算能力まとめ
HTC 10 Snapdragon 820 Kyro の浮動小数点演算能力
iPhone SE, Apple A9 の浮動小数点演算速度
ARM Cortex-A53 の浮動小数点演算速度とコンパイル時間の比較
iPod touch 6 の浮動小数点演算速度は Core 2 Duo ライン超え
iPad Air 2 (Apple A8X) の浮動小数点演算能力
ARM cpu vfp の種類と fp16 命令を使ってみる