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

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 命令を使ってみる