Waveglow_Inference_in_CUDAのROCm移植しようとした時の記録

とあるCUDA実装をROCｍに移植したのでその過程をメモに残します。
一ヶ月ほどに渡って奮闘しましたがあまり成果が出ませんでした
結論から言うとcudnn関連のAPI移植がどうすればいいかわからなかった点です

環境

OS Ubuntu 18.04LST
NVIDIA CUDA用GPU　RTX2080TI
CUDA version CUDA10.2
ROCm用GPU RadeonⅦ
ROCm vesion rocm-3.1.0

python環境はminicondaで作成（デフォルトはpython3.7です）

まずはCUDAで動かす

そもそもCUDA10.1で動かない問題があったので以下のプルリクを反映させた

Makefileを以下のように修正

DEBUG:= -D DEBUG

TARGET    = waveglow_tts
SRC_DIR   = waveglow/src
SRC_DIR_SYS   = sys/src
SRC_DIR_COMMON   = common/src

OBJ_DIR   = waveglow/obj
OBJ_DIR_SYS   = sys/obj
OBJ_DIR_COMMON   = common/obj

INCLUDES:=-Iwaveglow/header/ -Icommon/header -Isys/header -I/usr/local/cuda/include
NVCC:=nvcc
LDFLAGS:= -lcudnn -lcublas -lcurand
NVCCFLAGS:= -arch=sm_70 -std=c++11   #--compiler-bindir=g++-4.9 #--ptxas-options=-v
CUDNN_PATH:= /usr/local/cuda/
LIBS:= -L $(CUDNN_PATH)/lib64 -L/usr/local/lib

リポジトリのREADMEに従って waveglow_weights, mel_spectrogramsを解凍してローカルリポジトリの中に展開します。

後hparams.hppにwaveglow_weights pathを追加する必要があるので追加してください
mel_spectrogramsの中身は直接Waveglow_Inference_in_CUDAディレクトリにツッコんでfilenameに

LJ001-0015.wav.npy
LJ001-0051.wav.npy
LJ001-0063.wav.npy
LJ001-0072.wav.npy
LJ001-0079.wav.npy
LJ001-0094.wav.npy
LJ001-0096.wav.npy
LJ001-0102.wav.npy
LJ001-0153.wav.npy
LJ001-0173.wav.npy

となるようにして実行しました。

つぎにmakeします

$ make 
nvcc -arch=sm_70 -std=c++11    -lcudnn -lcublas -lcurand -L /usr/local/cuda//lib64 -L/usr/local/lib -o waveglow_tts common/obj/utils.o common/obj/logger.o common/obj/cnpy.o sys/obj/dropout.o sys/obj/conv_grad.o sys/obj/dense_mv.o sys/obj/conv.o sys/obj/dense.o waveglow/obj/WN.o waveglow/obj/main.o waveglow/obj/upsample.o

こうれで上手くビルドできるはず

./waveglow_tts filename.txt  ./出力するpath/
pip install scipy
python tools/npy_2_aud.py ./出力するpath

LJ001-0015.wavなど数点の英語の合成音声が生成されれば成功です。

これをROCｍに実装する際のことを考える

これはCUDAで書かれていますがAMD ROCｍでは各種CUDA APIを互換APIでサポートしています

今回はcudnn関連のAPIが多いので参照してみます.
https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/CUDNN_API_supported_by_HIP.md

例えば
cudnnGetConvolutionBackwardDataWorkspaceSize()　であれば hipdnnGetConvolutionBackwardDataWorkspaceSize()　に置き換えることで互換性が確保されます.

ただ全部のcudnn関数がサポートされているわけではないので注意が必要です。

hipify-perlでコードを変更

hipify-perlで以下のコードを変更する
.cuは.cppにリネーム

common/build_memo
common/header/cnpy.hpp
common/header/data_types.hpp
common/header/logger.hpp
common/header/macros.hpp
common/header/utils.hpp
common/src/logger.cpp
sys/header/activation.hpp
sys/header/conv.hpp
sys/header/conv_grad.hpp
sys/header/conv_old.hpp
sys/header/dense.hpp
sys/header/dense_mv.hpp
sys/header/dropout.hpp
sys/src/conv.cpp
sys/src/conv_grad.cpp
sys/src/dense.cpp
sys/src/dense_mv.cpp
sys/src/dropout.cpp
waveglow/header/WN.hpp
waveglow/header/hparams.hpp
waveglow/header/upsample.hpp
waveglow/src/WN.cpp
waveglow/src/main.cpp
waveglow/src/upsample.cpp

hipily-perlが追加してきた.hのビルド

cudnnの置換にhipdnn.hをインクルードしてきたのでhipdnnをビルドしてinstallした

hipdnnのinstall手順

git clone https://github.com/ROCmSoftwarePlatform/hipDNN.git
cd hipDNN/
mkdir build
cd build
cmake ..make
make
sudo make install 

hipdnn.hのリネーム

installされたincluidファイルはhipdnn.hですがhipily-perlはhipDNN.hと出力してくるのでhipdnn.hをcpコマンド絵

cd /opt/rocm/hipdnn/
sudo cp hipdnn.h hipDNN.h 

makefileを覗いてみてどのコードがビルドされるか観察する

.cuから.cppに全部変更してひとまずここまで変えました

DEBUG:= -D DEBUG

TARGET    = waveglow_tts
SRC_DIR   = waveglow/src
SRC_DIR_SYS   = sys/src
SRC_DIR_COMMON   = common/src

OBJ_DIR   = waveglow/obj
OBJ_DIR_SYS   = sys/obj
OBJ_DIR_COMMON   = common/obj

INCLUDES:=-Iwaveglow/header/ -Icommon/header -Isys/header  -I/opt/rocm-3.1.0/include/ -I/opt/rocm/hipdnn/include -I/opt/rocm/hiprand/include -I/opt/rocm-3.1.0/rocrand/include #-I/usr/local/cuda/include

NVCC:=hipcc
LDFLAGS:=   -lcudnn -lcublas -lcurand 
NVCCFLAGS:=  -std=c++11   #-amdgpu-target=gfx906 #--compiler-bindir=g++-4.9 #--ptxas-options=-v -arch=sm_70
CUDNN_PATH:= /usr/local/cuda/
LIBS:= -L $(CUDNN_PATH)/lib64 -L/usr/local/lib  -L/opt/rocm-3.1.0/lib/  -L/opt/rocm/lib   -L/opt/rocm/hipdnn/lib -L/opt/rocm-3.1.0/rocrand/lib -L/opt/rocm/hipdnn/lib
# -L/opt/rocm-3.1.0/hsa/lib/ 



CU_FILES  = $(wildcard $(SRC_DIR)/*.cpp)
OBJS = $(patsubst %.cpp,$(OBJ_DIR)/%.o,$(notdir $(CU_FILES)))

CU_FILES_SYS  = $(wildcard $(SRC_DIR_SYS)/*.cpp)
OBJS_SYS = $(patsubst %.cpp,$(OBJ_DIR_SYS)/%.o,$(notdir $(CU_FILES_SYS)))

CU_FILES_COMMON  = $(wildcard $(SRC_DIR_COMMON)/*.cpp)
OBJS_COMMON = $(patsubst %.cpp,$(OBJ_DIR_COMMON)/%.o,$(notdir $(CU_FILES_COMMON)))


$(TARGET) :	dirmake $(OBJS_COMMON)	$(OBJS_SYS)	$(OBJS)
	$(NVCC) $(NVCCFLAGS) $(LDFLAGS) $(LIBS) -o $@ $(OBJS_COMMON) $(OBJS_SYS) $(OBJS)


$(OBJ_DIR)/%.o : $(SRC_DIR)/%.cpp
	$(NVCC) $(NVCCFLAGS) $(INCLUDES) -o $@ -c $< $(DEBUG)

$(OBJ_DIR_SYS)/%.o : $(SRC_DIR_SYS)/%.cpp
	$(NVCC) $(NVCCFLAGS) $(INCLUDES) -o $@ -c $< $(DEBUG)

$(OBJ_DIR_COMMON)/%.o : $(SRC_DIR_COMMON)/%.cpp
	$(NVCC) $(NVCCFLAGS) $(INCLUDES) -o $@ -c $< $(DEBUG)


dirmake:
	@mkdir -p $(OBJ_DIR_COMMON)
	@mkdir -p $(OBJ_DIR_SYS)
	@mkdir -p $(OBJ_DIR)

.PHONY :	clean
clean :	
	rm -f $(TARGET)
	rm -f $(OBJ_DIR_COMMON)/*.o
	rm -f $(OBJ_DIR_SYS)/*.o
	rm -f $(OBJ_DIR)/*.o

rebuild:	clean build

主なmake時のエラーなど

この時点は500行ほど出力されます。

HIPDNN_TENSOR_OP_MATH_ALLOW_CONVERSION

CUDNNでtensorcoreを使うときに明示するものをHIPで置き換えたものとおもわれます
hipilyでは以下のような置換が実装されていますが

{"CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION",               {"HIPDNN_TENSOR_OP_MATH_ALLOW_CONVERSION",               "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},    // 2

宣言が見つからないとコンパイラに怒られてしまいました

sys/src/conv.cpp:117:66: error: use of undeclared identifier 'HIPDNN_TENSOR_OP_MATH_ALLOW_CONVERSION'

hipdnnSetFilter4dDescriptor

hipdnnSetFilter4dDescriptorはhipdnnでサポートが追加されました。

hipdnnのinclude pathは通しましたがエラーの状態です

sys/src/conv_grad.cpp:54:24: error: no matching function for call to 'hipdnnSetFilter4dDescriptor'
            checkCUDNN(hipdnnSetFilter4dDescriptor(kernel_descriptor,

正しく関数が呼び出せてないとエラーがでている。

統括

cudnn⇨hipdnnへの移植が難しくエラーを取り切れなかった

hipdnnの更新は2019年で止まっておりDocumentも少ない為対応が難しい。

rocmのpathについて

rocm3.0.1から/opt/rocmから/opt/rocm{version}/にrocmディレクトリの標準が変更になり古いrocm関連のツールではopt/rocm/にinstallされることがあるため互換性の点で厳しい気がします。。