[APP] ROCM RPC (#1155)

apps/rocm_rpc/Makefile

+# Makefile Example to deploy TVM modules.
+ROCM_PATH=/opt/rocm
+
+TVM_ROOT=$(shell cd ../..; pwd)
+NNVM_PATH=nnvm
+DMLC_CORE=${TVM_ROOT}/dmlc-core
+
+PKG_CFLAGS = -std=c++11 -O2 -fPIC\
+	-I${TVM_ROOT}/include\
+	-I${DMLC_CORE}/include\
+	-I${TVM_ROOT}/dlpack/include\
+	-I${ROCM_PATH}/include
+
+PKG_LDFLAGS = -L${ROCM_PATH}/lib -L${TVM_ROOT}/lib -ldl -lpthread -lhip_hcc -lMIOpen
+
+.PHONY: clean all
+
+all: lib/libtvm_runtime_rocm.so
+
+# Build rule for all in one TVM package library
+lib/libtvm_runtime_rocm.so: rocm_runtime_pack.cc
+	@mkdir -p $(@D)
+	$(CXX) $(PKG_CFLAGS) -shared -o $@ $(filter %.cc %.o %.a, $^) $(PKG_LDFLAGS)

apps/rocm_rpc/README.md

+# TVM ROCm RPC
+
+This folder contains a simple recipe to make RPC work together with ROCm.TVM's RPC server relies on process
+fork to create a new process for each incoming session.
+Like CUDA, opencl driver, the runtime ROCm runtime is not fork-safe.
+A typical CUDA or opencl driver will initialize lazily
+and we can use normal TVM RPC server because we won't touch the driver API before we fork a new session.
+However, the current ROCm runtime eagerly initialize during startup and will directly cause error during fork.
+This folder provides a workaround to this problem.
+
+## Usage
+- Build tvm **without** rocm (it is important to exclude rocm from runtime)
+- Modify the ROCM_PATH to be the correct path the current [Makefile](Makefile)
+- Type make to build lib/libtvm_runtime_rocm.so, which is a standalone dll module
+- Use [start_rpc_server.sh](start_rpc_server.sh) to start the RPC server
+
+## How it works
+- The RPC server starts without ROCm dependency.
+- lib/libtvm_runtim_rocm.so is dynamically loaded only after the fork.
+
+## Note
+With ROCm RPC, we can build AMDGPU program from a machine without AMD GPU
+and remotely upload and execute on a AMDGPU machine.
+Please note that you will need to set the gfx version correctly(via ```-model```` or ```-mcpu```)
+because we can no longer query the GPU version dynamically during runtime.
+
+
+```python
+import tvm
+from tvm.contrib import rpc
+
+# set mcpu explicitly to be the gpu version.
+target = "rocm -mcpu=gfx900"
+remote = rpc.connect(server_host, server_port)
+mod = tvm.build(s, args, target)
+mod.export_library("mylib.so")
+
+remote.upload("mylib.so")
+foo = remote.load_module("mylib.so")
+# same as normal RPC
+```

apps/rocm_rpc/rocm_runtime_pack.cc

+/*!
+ * \brief This is an all in one file for ROCM runtime library.
+ *
+ * This is used to create a RPC module library that can be
+ * safely passed to rocm
+ */
+
+#define TVM_ROCM_RUNTIME 1
+#define TVM_USE_MIOPEN 1
+#define __HIP_PLATFORM_HCC__ 1
+
+#include "../../src/runtime/rocm/rocm_device_api.cc"
+#include "../../src/runtime/rocm/rocm_module.cc"
+#include "../../src/contrib/miopen/conv_forward.cc"
+#include "../../src/contrib/miopen/miopen_utils.cc"

apps/rocm_rpc/start_rpc_server.sh

+#/bin/bash
+PROJ_ROOT=$(realpath $(dirname "$0")/../..)
+export PYTHONPATH=${PROJ_ROOT}/python:${PYTHONPATH}
+
+python -m tvm.exec.rpc_server "$@" --load-library=${PROJ_ROOT}/apps/rocm_rpc/lib/libtvm_runtime_rocm.so

python/tvm/exec/rpc_server.py

@@ -2,6 +2,7 @@
 from __future__ import absolute_import
 
 import argparse
+import logging
 from ..contrib import rpc
 
 def main():
@@ -31,6 +32,7 @@ def main():
     else:
         tracker_addr = None
 
+    logging.basicConfig(level=logging.INFO)
     server = rpc.Server(args.host,
                         args.port,
                         args.port_end,