diff --git a/CMakeLists.txt b/CMakeLists.txt
index abf9f4a2a945a..2a4dfbafd2449 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -30,6 +30,7 @@ tvm_option(USE_ROCM "Build with ROCM" OFF)
 tvm_option(ROCM_PATH "The path to rocm" /opt/rocm)
 tvm_option(USE_HEXAGON_DEVICE "Build with Hexagon device support in TVM runtime" OFF)
 tvm_option(USE_HEXAGON_SDK "Path to the Hexagon SDK root (required for Hexagon support in TVM runtime or for building TVM runtime for Hexagon)" /path/to/sdk)
+tvm_option(USE_HEXAGON_RPC "Enable Hexagon RPC using minRPC implementation over Android." OFF)
 tvm_option(USE_HEXAGON_LAUNCHER "Build the Hexagon graph launcher application" OFF)
 tvm_option(USE_HEXAGON_PROXY_RPC "Build the Hexagon Proxy RPC server application" OFF)
 tvm_option(USE_RPC "Build with RPC" ON)
diff --git a/apps/cpp_rpc/CMakeLists.txt b/apps/cpp_rpc/CMakeLists.txt
index ccac53fc3ca04..966448929a0ac 100644
--- a/apps/cpp_rpc/CMakeLists.txt
+++ b/apps/cpp_rpc/CMakeLists.txt
@@ -6,6 +6,8 @@ set(TVM_RPC_SOURCES
   rpc_server.cc
 )
 
+set(TVM_RPC_LINKER_LIBS "")
+
 if(WIN32)
   list(APPEND TVM_RPC_SOURCES win32_process.cc)
 endif()
@@ -43,4 +45,12 @@ target_include_directories(
   PUBLIC DMLC_PATH
 )
 
-target_link_libraries(tvm_rpc tvm_runtime)
+if (BUILD_FOR_ANDROID AND USE_HEXAGON_SDK)
+  find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}")
+  link_directories(${HEXAGON_REMOTE_ROOT})
+  list(APPEND TVM_RPC_LINKER_LIBS cdsprpc log)
+endif()
+
+list(APPEND TVM_RPC_LINKER_LIBS tvm_runtime)
+ 
+target_link_libraries(tvm_rpc ${TVM_RPC_LINKER_LIBS})
diff --git a/cmake/libs/hexagon_rpc_skel/CMakeLists.txt b/cmake/libs/hexagon_rpc_skel/CMakeLists.txt
new file mode 100644
index 0000000000000..a4756aa4e0881
--- /dev/null
+++ b/cmake/libs/hexagon_rpc_skel/CMakeLists.txt
@@ -0,0 +1,119 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License.  You may obtain a copy of the License at
+#
+#   http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied.  See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+cmake_minimum_required(VERSION 3.2)
+include(ExternalProject)
+project(HexagonRPCSkel C CXX)
+
+set(TVM_SOURCE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/../../..")
+set(TVM_SRC_DIR "${TVM_SOURCE_DIR}/src")
+
+
+include("${TVM_SOURCE_DIR}/cmake/modules/HexagonSDK.cmake")
+
+find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}")
+
+include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_REMOTE_ROOT})
+
+set(HEXAGON_RPC_H "hexagon_rpc.h")
+set(HEXAGON_RPC_SKEL_C "hexagon_rpc_skel.c")
+set(HEXAGON_RPC_STUB_C "hexagon_rpc_stub.c")
+
+include_directories(
+  "${TVM_SOURCE_DIR}/include"
+  "${TVM_SOURCE_DIR}/3rdparty/dlpack/include"
+  "${TVM_SOURCE_DIR}/3rdparty/dmlc-core/include"
+)
+
+set(QAIC_EXE "${HEXAGON_QAIC_EXE}")
+foreach(INCDIR IN LISTS HEXAGON_SDK_INCLUDES HEXAGON_REMOTE_ROOT)
+  list(APPEND QAIC_FLAGS "-I${INCDIR}")
+endforeach()
+
+add_custom_command(
+  OUTPUT ${HEXAGON_RPC_SKEL_C} ${HEXAGON_RPC_H}
+  COMMAND ${QAIC_EXE} ${QAIC_FLAGS} "${TVM_SRC_DIR}/runtime/hexagon/rpc/hexagon_rpc.idl"
+  MAIN_DEPENDENCY "${TVM_SRC_DIR}/runtime/hexagon/rpc/hexagon_rpc.idl"
+)
+
+include_directories(SYSTEM
+  ${HEXAGON_QURT_INCLUDES}
+  ${CMAKE_CURRENT_BINARY_DIR}   # Output of qaic will go here
+)
+
+link_directories(${HEXAGON_QURT_LIBS})
+
+add_definitions(-D_MACH_I32=int)
+add_definitions(-DDMLC_CXX11_THREAD_LOCAL=0)
+add_definitions(-DDMLC_USE_LOGGING_LIBRARY=<tvm/runtime/logging.h>)
+
+# Extra compile flags (both C and C++).
+set(EXTRA_COMP_FLAGS
+  "-O3"
+  "-m${USE_HEXAGON_ARCH}"
+)
+string(REGEX REPLACE ";" " " EXTRA_COMP_FLAGS_STR "${EXTRA_COMP_FLAGS}")
+set(CMAKE_C_FLAGS "${EXTRA_COMP_FLAGS_STR} ${CMAKE_C_FLAGS}")
+set(CMAKE_CXX_FLAGS "${EXTRA_COMP_FLAGS_STR} ${CMAKE_CXX_FLAGS}")
+
+set(SKEL_SRCS
+  "${TVM_SRC_DIR}/runtime/hexagon/rpc/hexagon/rpc_server.cc"
+)
+
+set(MINRPC_SRCS
+  "${TVM_SRC_DIR}/runtime/minrpc/minrpc_server.h"
+  "${TVM_SRC_DIR}/runtime/minrpc/rpc_reference.h"
+)
+
+set(TVM_RPC_SRC 
+  "${TVM_SRC_DIR}/runtime/rpc/rpc_module.cc"
+  "${TVM_SRC_DIR}/runtime/rpc/rpc_endpoint.cc"
+  "${TVM_SRC_DIR}/runtime/rpc/rpc_session.cc"
+  "${TVM_SRC_DIR}/runtime/rpc/rpc_local_session.cc"
+)
+
+add_library(hexagon_rpc_skel SHARED
+  "${HEXAGON_RPC_H}"
+  "${HEXAGON_RPC_SKEL_C}"
+  "${SKEL_SRCS}"
+  "${MINRPC_SRCS}"
+  "${TVM_RPC_SRC}"
+)
+
+ExternalProject_Add(static_hexagon_tvm_runtime
+  SOURCE_DIR "${TVM_SOURCE_DIR}"
+  BUILD_COMMAND $(MAKE) runtime
+  CMAKE_ARGS
+  "-DCMAKE_C_COMPILER=${CMAKE_C_COMPILER}"
+  "-DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER}"
+  "-DUSE_HEXAGON_ARCH=${USE_HEXAGON_ARCH}"
+  "-DCMAKE_CXX_STANDARD=14"
+  "-DUSE_LIBBACKTRACE=OFF"
+  "-DUSE_LLVM=OFF"
+  "-DUSE_RPC=OFF"
+  "-DBUILD_STATIC_RUNTIME=ON"
+  "-DUSE_HEXAGON_SDK=${USE_HEXAGON_SDK}"
+  INSTALL_COMMAND ""
+  BUILD_ALWAYS ON
+)
+ExternalProject_Get_Property(static_hexagon_tvm_runtime BINARY_DIR)
+
+add_dependencies(hexagon_rpc_skel static_hexagon_tvm_runtime)
+add_library(h_tvm_runtime STATIC IMPORTED)
+set_target_properties(h_tvm_runtime PROPERTIES IMPORTED_LOCATION "${BINARY_DIR}/libtvm_runtime.a")
+
+target_link_libraries(hexagon_rpc_skel -Wl,--whole-archive h_tvm_runtime -Wl,--no-whole-archive)
diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake
index 4c12bedde2af4..43789dda5c45c 100644
--- a/cmake/modules/Hexagon.cmake
+++ b/cmake/modules/Hexagon.cmake
@@ -53,7 +53,6 @@ if(BUILD_FOR_HEXAGON)
   include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_QURT_INCLUDES})
 endif()
 
-
 if (NOT USE_HEXAGON_SDK STREQUAL "" AND
     NOT USE_HEXAGON_SDK STREQUAL "/path/to/sdk")
   set(HEXAGON_SDK_PATH_DEFINED ${USE_HEXAGON_SDK})
@@ -73,10 +72,12 @@ endif()
 # e.g. when compiling the TVM runtime for Hexagon.
 if (NOT BUILD_FOR_HEXAGON AND NOT BUILD_FOR_ANDROID)
   if(USE_HEXAGON_LAUNCHER STREQUAL "OFF" AND
-      USE_HEXAGON_PROXY_RPC STREQUAL "OFF")
+      USE_HEXAGON_PROXY_RPC STREQUAL "OFF" AND NOT USE_HEXAGON_RPC)
     if(USE_HEXAGON_DEVICE STREQUAL "OFF")
       list(APPEND COMPILER_SRCS src/target/opt/build_hexagon_off.cc)
-      return()
+      if (NOT USE_HEXAGON_RPC)
+        return()
+      endif()
     elseif(NOT USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}" AND
         NOT USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}")
       set(ERROR_MSG
@@ -202,6 +203,103 @@ if(USE_HEXAGON_PROXY_RPC STREQUAL "ON")
   set_directory_properties(PROPERTIES ADDITIONAL_MAKE_CLEAN_FILES "${RPC_BINARY_DIR}")
 endif()
 
+if(USE_HEXAGON_RPC)
+  if(DEFINED USE_ANDROID_TOOLCHAIN)
+    if(NOT DEFINED ANDROID_PLATFORM)
+      message(SEND_ERROR "Please set ANDROID_PLATFORM "
+        "when providing an Android cmake toolchain.")
+    endif()
+    if(NOT DEFINED ANDROID_ABI)
+      message(SEND_ERROR "Please set ANDROID_ABI "
+        "when providing an Android cmake toolchain.")
+    endif()
+  else()
+    message(SEND_ERROR "Please set USE_ANDROID_TOOLCHAIN to build the android "
+      "RPC server for Hexagon.")
+  endif()
+
+  if(NOT DEFINED USE_HEXAGON_SDK)
+    message(SEND_ERROR "Please set USE_HEXAGON_SDK to build the android "
+      "RPC server for Hexagon RPC.")
+  endif()
+  if(NOT DEFINED USE_HEXAGON_ARCH)
+    message(SEND_ERROR "Please set USE_HEXAGON_ARCH to build the android "
+      "RPC server for Hexagon RPC.")
+  endif()
+  find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}")
+
+  set(HEXAGON_RPC_OUTPUT "${CMAKE_CURRENT_BINARY_DIR}/hexagon_rpc")
+  file(MAKE_DIRECTORY ${HEXAGON_RPC_OUTPUT})
+
+  # Android Part
+  ExternalProject_Add(android_runtime_rpc
+    SOURCE_DIR "${CMAKE_SOURCE_DIR}"
+    BUILD_COMMAND $(MAKE) runtime tvm_rpc
+    CMAKE_ARGS
+    "-DCMAKE_TOOLCHAIN_FILE=${USE_ANDROID_TOOLCHAIN}"
+    "-DUSE_ANDROID_TOOLCHAIN=${USE_ANDROID_TOOLCHAIN}"
+    "-DANDROID_PLATFORM=${ANDROID_PLATFORM}"
+    "-DANDROID_ABI=${ANDROID_ABI}"
+    "-DCMAKE_CXX_STANDARD=14"
+    "-DUSE_LIBBACKTRACE=OFF"
+    "-DUSE_LLVM=OFF"
+    "-DUSE_RPC=ON"
+    "-DUSE_CPP_RPC=ON"
+    "-DUSE_HEXAGON_SDK=${USE_HEXAGON_SDK}"
+    "-DUSE_HEXAGON_ARCH=${USE_HEXAGON_ARCH}"
+    "-DCMAKE_VERBOSE_MAKEFILE=ON"
+    INSTALL_COMMAND ""
+    BUILD_ALWAYS ON
+  )
+  ExternalProject_Get_Property(android_runtime_rpc BINARY_DIR)
+  ExternalProject_Add_Step(android_runtime_rpc copy_binary_runtime
+    COMMAND ${CMAKE_COMMAND} -E copy_if_different
+      ${BINARY_DIR}/libtvm_runtime.so
+      ${HEXAGON_RPC_OUTPUT}/libtvm_runtime.so
+    DEPENDEES install
+  )
+  ExternalProject_Add_Step(android_runtime_rpc copy_binary_rpc
+    COMMAND ${CMAKE_COMMAND} -E copy_if_different
+      ${BINARY_DIR}/tvm_rpc
+      ${HEXAGON_RPC_OUTPUT}/tvm_rpc_android
+    DEPENDEES install
+  )
+
+  if("${USE_HEXAGON_TOOLCHAIN}" STREQUAL "")
+    message(SEND_ERROR "Please set USE_HEXAGON_TOOLCHAIN to build the hexagon "
+      "RPC SKEL.")
+  endif()
+  find_hexagon_toolchain()
+  message(STATUS "HEXAGON_TOOLCHAIN: ${HEXAGON_TOOLCHAIN}")
+
+  # Hexagon Part
+  ExternalProject_Add(hexagon_rpc_skel
+    SOURCE_DIR "${CMAKE_SOURCE_DIR}/cmake/libs/hexagon_rpc_skel"
+    INSTALL_DIR "${LAUNCHER_BINARY_DIR}"
+    CMAKE_ARGS
+    "-DCMAKE_C_COMPILER=${HEXAGON_TOOLCHAIN}/bin/hexagon-clang"
+    "-DCMAKE_CXX_COMPILER=${HEXAGON_TOOLCHAIN}/bin/hexagon-clang++"
+    "-DFASTRPC_LIBS=SKEL"
+    "-DUSE_HEXAGON_ARCH=${USE_HEXAGON_ARCH}"
+    "-DUSE_HEXAGON_SDK=${USE_HEXAGON_SDK}"
+    INSTALL_COMMAND ""
+    BUILD_ALWAYS ON
+  )
+  ExternalProject_Get_Property(hexagon_rpc_skel BINARY_DIR)
+  ExternalProject_Add_Step(hexagon_rpc_skel copy_hexagon_skel
+    COMMAND ${CMAKE_COMMAND} -E copy_if_different
+      ${BINARY_DIR}/libhexagon_rpc_skel.so
+      ${HEXAGON_RPC_OUTPUT}/libhexagon_rpc_skel.so
+    DEPENDEES install
+  )
+
+  # copy android_bash template file
+  configure_file("${CMAKE_SOURCE_DIR}/src/runtime/hexagon/rpc/android_bash.sh.template" 
+    ${HEXAGON_RPC_OUTPUT} COPYONLY)
+
+  set_directory_properties(PROPERTIES ADDITIONAL_MAKE_CLEAN_FILES "${HEXAGON_RPC_OUTPUT}")
+endif()
+
 if(USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}")
   find_hexagon_toolchain()
   message(STATUS "Hexagon toolchain: ${HEXAGON_TOOLCHAIN}")
@@ -227,6 +325,7 @@ elseif(USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}")
     ${HEXAGON_RPCMEM_ROOT}/inc
     ${HEXAGON_REMOTE_ROOT}
   )
+
   list(APPEND TVM_RUNTIME_LINKER_LIBS "dl")
   if(BUILD_FOR_ANDROID)
     # Hexagon runtime uses __android_log_print, which is in liblog.
@@ -241,10 +340,39 @@ if (USE_HEXAGON_DEVICE STREQUAL "${PICK_NONE}")
   elseif(BUILD_FOR_ANDROID AND HEXAGON_SDK_PATH_DEFINED)
     list(APPEND RUNTIME_HEXAGON_SRCS src/runtime/hexagon/proxy_rpc/device_api.cc)
   else()
-    file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/host/*.cc)
+  file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/host/*.cc)
   endif()
 else()
   file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/android/*.cc)
 endif()
+
+if(USE_HEXAGON_RPC)
+  file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/host/*.cc)
+endif()
+
+if(USE_HEXAGON_SDK AND BUILD_FOR_ANDROID)
+  find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}")
+  include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_REMOTE_ROOT})
+
+  set(QAIC_EXE "${HEXAGON_QAIC_EXE}")
+  foreach(INCDIR IN LISTS HEXAGON_SDK_INCLUDES HEXAGON_REMOTE_ROOT)
+    list(APPEND QAIC_FLAGS "-I${INCDIR}")
+  endforeach()
+
+  set(HEXAGON_RPC_DIR "${CMAKE_SOURCE_DIR}/src/runtime/hexagon/rpc")
+  set(RPC_IDL "hexagon_rpc.idl")
+  set(RPC_H "hexagon_rpc.h")
+  set(RPC_STUB_C "hexagon_rpc_stub.c")
+  
+  add_custom_command(
+    OUTPUT "${HEXAGON_RPC_DIR}/${RPC_STUB_C}" "${HEXAGON_RPC_DIR}/${RPC_H}"
+    COMMAND ${QAIC_EXE} ${QAIC_FLAGS} "${HEXAGON_RPC_DIR}/${RPC_IDL}" -o ${HEXAGON_RPC_DIR}
+    MAIN_DEPENDENCY "${HEXAGON_RPC_DIR}/${RPC_IDL}"
+  )
+  file(GLOB HEXAGON_RPC_CPP "${HEXAGON_RPC_DIR}/android/*.cc")
+  set(HEXAGON_RPC_STUB_C "${HEXAGON_RPC_DIR}/${RPC_STUB_C}")
+endif()
+
 list(APPEND RUNTIME_SRCS ${RUNTIME_HEXAGON_SRCS} ${RUNTIME_HEXAGON_SIM_SRCS}
-                         ${RUNTIME_HEXAGON_DEVICE_SRCS} ${RUNTIME_HEXAGON_COMMON_SRCS})
+                         ${RUNTIME_HEXAGON_DEVICE_SRCS} ${HEXAGON_RPC_CPP} ${HEXAGON_RPC_STUB_C} 
+                         ${RUNTIME_HEXAGON_COMMON_SRCS})
diff --git a/python/tvm/contrib/hexagon/__init__.py b/python/tvm/contrib/hexagon/__init__.py
new file mode 100644
index 0000000000000..e728216853e5e
--- /dev/null
+++ b/python/tvm/contrib/hexagon/__init__.py
@@ -0,0 +1,17 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License.  You may obtain a copy of the License at
+#
+#   http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied.  See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Hexagon APIs."""
diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py
new file mode 100644
index 0000000000000..ef081f2d79e3e
--- /dev/null
+++ b/python/tvm/contrib/hexagon/build.py
@@ -0,0 +1,300 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License.  You may obtain a copy of the License at
+#
+#   http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied.  See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+"""Defines top-level glue functions for building Hexagon."""
+
+import pathlib
+import os
+import subprocess
+from typing import Union
+import stat
+import datetime
+
+import tvm
+from ..._ffi import libinfo
+from .session import Session
+
+
+RPC_SERVER_FILES = ["tvm_rpc_android", "libtvm_runtime.so", "android_bash.sh"]
+
+HEXAGON_FILES = ["libhexagon_rpc_skel.so"]
+
+HEXAGON_RPC_DIR = None
+
+ANDROID_HEXAGON_TEST_BASE_DIR = pathlib.Path("/data/local/tmp/hexagon_test")
+
+
+def get_hexagon_rpc_dir() -> pathlib.Path:
+    """Find the Hexagon library.
+
+    Returns
+    -------
+    str :
+        The path to the Hexagon library
+    """
+    global HEXAGON_RPC_DIR
+    if HEXAGON_RPC_DIR is None:
+        for path in libinfo.find_lib_path():
+            rpc_dir = os.path.join(os.path.dirname(path), "hexagon_rpc")
+            if os.path.isdir(rpc_dir):
+                HEXAGON_RPC_DIR = pathlib.Path(rpc_dir)
+                break
+        else:
+            raise "hexagon_rpc was not found."
+    return HEXAGON_RPC_DIR
+
+
+class HexagonLauncher:
+    """Hexagon Launcher"""
+
+    def __init__(self, serial_number: str, workspace_size_gb: int = 1):
+        """Configure a new HexagonLauncher
+
+        Parameters
+        ----------
+        serial_number : str
+            Android device serial number from android 'adb' command.
+        """
+        # Hexagon RPCSession
+        self.session = None
+
+        self._serial_number = serial_number
+        self._adb_device_sub_cmd = ["adb", "-s", self._serial_number]
+        self._mod = None
+        self._workspace = None
+        self._workspace_max_size_mb = workspace_size_gb * 1024
+
+    HEXAGON_REMOTE_DEVICE_KEY = "hexagon-dev"
+
+    def android_run_rpc(
+        self,
+        workspace_dir: Union[str, pathlib.Path] = None,
+        rpc_server_port: int = 7070,
+        rpc_tracker_host: str = "0.0.0.0",
+        rpc_tracker_port: int = 9190,
+    ):
+        """Upload Android artifacts and run RPC server on Android.
+
+        Parameters
+        ----------
+        workspace_dir : Union[str, pathlib.Path]
+            Workspace directory used on Android to upload artifacts.
+
+        rpc_server_port : int
+            Android RPC server port number
+
+        rpc_tracker_host : str
+            RPC tracker IP on host
+
+        rpc_tracker_port : int
+            RPC tracker port on host
+        """
+        # Create test base directory
+        subprocess.check_call(
+            self._adb_device_sub_cmd + ["shell", "mkdir", "-p", ANDROID_HEXAGON_TEST_BASE_DIR]
+        )
+
+        # Check size of base directory and cleanup if needed
+        while self._get_workspace_size() > self._workspace_max_size_mb:
+            self._workspace_remove_latest()
+
+        if not workspace_dir:
+            self._workspace = str(
+                ANDROID_HEXAGON_TEST_BASE_DIR
+                / datetime.datetime.now().strftime("%Y-%m-%dT%H-%M-%S")
+            )
+        else:
+            self._workspace = workspace_dir
+
+        # Upload RPC server and libraries
+        subprocess.check_call(self._adb_device_sub_cmd + ["shell", "mkdir", "-p", self._workspace])
+
+        # Create bash script
+        android_bash_script_path = get_hexagon_rpc_dir() / "android_bash.sh"
+        with open(get_hexagon_rpc_dir() / "android_bash.sh.template", "r") as src_f:
+            if os.path.exists(android_bash_script_path):
+                os.remove(android_bash_script_path)
+            with open(android_bash_script_path, "w") as dest_f:
+                for line in src_f.readlines():
+                    if "<RPC_TRACKER_HOST>" in line:
+                        line = line.replace("<RPC_TRACKER_HOST>", str(rpc_tracker_host))
+                    if "<RPC_TRACKER_PORT>" in line:
+                        line = line.replace("<RPC_TRACKER_PORT>", str(rpc_tracker_port))
+                    if "<HEXAGON_REMOTE_DEVICE_KEY>" in line:
+                        line = line.replace(
+                            "<HEXAGON_REMOTE_DEVICE_KEY>", self.HEXAGON_REMOTE_DEVICE_KEY
+                        )
+                    if "<RPC_SERVER_PORT>" in line:
+                        line = line.replace("<RPC_SERVER_PORT>", str(rpc_server_port))
+                    dest_f.write(line)
+
+        # Make shell script executable
+        android_bash_stat = os.stat(android_bash_script_path)
+        os.chmod(android_bash_script_path, android_bash_stat.st_mode | stat.S_IEXEC)
+
+        # Push files
+        for item in RPC_SERVER_FILES:
+            src_path = get_hexagon_rpc_dir() / item
+            destination = f"{self._workspace}/{item}"
+            subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, destination])
+
+        # Removed pre-defined forward/reverse rules
+        subprocess.check_call(self._adb_device_sub_cmd + ["forward", "--remove-all"])
+        subprocess.check_call(self._adb_device_sub_cmd + ["reverse", "--remove-all"])
+
+        # Enable port reverse for RPC tracker
+        subprocess.check_call(
+            self._adb_device_sub_cmd
+            + ["reverse", f"tcp:{rpc_tracker_port}", f"tcp:{rpc_tracker_port}"]
+        )
+        # Enable port forward for RPC server. We forward 9 ports after the rpc_server_port.
+        for i in range(0, 10):
+            subprocess.check_call(
+                self._adb_device_sub_cmd
+                + ["forward", f"tcp:{rpc_server_port+i}", f"tcp:{rpc_server_port+i}"]
+            )
+
+        # Run server and connect to tracker
+        subprocess.Popen(
+            self._adb_device_sub_cmd + ["shell", f"cd {self._workspace} && ./android_bash.sh"],
+            stdout=subprocess.PIPE,
+            stdin=subprocess.PIPE,
+            stderr=subprocess.PIPE,
+        )
+
+    def hexagon_setup(self):
+        """Upload Hexagon artifacts on Android."""
+        for item in HEXAGON_FILES:
+            src_path = get_hexagon_rpc_dir() / item
+            dst_path = f"{self._workspace}/{item}"
+            subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, dst_path])
+
+    def hexagon_session_setup(self, remote_kw: dict):
+        """Setup Hexagon RPC Session from host to Hexagon device.
+
+        Parameters
+        ----------
+        remote_kw : dict
+            RPC tracker configs.
+        """
+        hexagon_remote_kw = dict(remote_kw)
+        hexagon_remote_kw["key"] = self.HEXAGON_REMOTE_DEVICE_KEY
+        self.session = Session(hexagon_remote_kw)
+
+    def get_module(self, module_name: str):
+        """Load a Hexagon TVM module, already uploaded on Android, on Hexagon and return the module.
+
+        Parameters
+        ----------
+        module_name : str
+            Module filename.
+
+        Returns
+        -------
+        TVMModule :
+            A TVM Module loaded on hexagon.
+        """
+        module_path = f"{self._workspace}/{module_name}"
+        self._mod = self.session.load_module(module_path)
+        return self._mod
+
+    def upload(self, host_path: Union[str, pathlib.Path], remote_filename: str):
+        """Upload a file to remote(Android).
+
+        Parameters
+        ----------
+        host_path : Union[str, pathlib.Path]
+            File path on host.
+
+        remote_filename : str
+            File name on remote(Android).
+        Returns
+        -------
+        TVMModule :
+            A TVM Module loaded on hexagon.
+        """
+        src_path = str(host_path)
+        dst_remote_path = f"{self._workspace}/{remote_filename}"
+        subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, dst_remote_path])
+
+    def get_graph_executor(self, libmod, remote_libmod_filename: str):
+        """Create a local GraphModule which consumes a remote libmod.
+
+        Parameters
+        ----------
+        libmod : tvm.runtime.Module
+            The module of the corresponding function.
+            This library module is for remote hexagon runtime.
+
+        remote_libmod_filename : str
+            Module filename on remote. It is assumed this file lives under self._workspace path.
+
+        Returns
+        -------
+        graph_module : GraphModule
+            Runtime graph module that can be used to execute the graph.
+        """
+        self.session.__enter__()
+        hexagon_mod = self.get_module(remote_libmod_filename)
+        return tvm.contrib.graph_executor.create(
+            libmod.get_graph_json(), hexagon_mod, self.session.device
+        )
+
+    def close(self):
+        """Close RPC server on Android"""
+        # Kill process childs
+        subprocess.Popen(
+            self._adb_device_sub_cmd + ["shell", f"pkill -P `cat {self._workspace}/rpc_pid.txt`"]
+        )
+        # Kill main process
+        subprocess.Popen(
+            self._adb_device_sub_cmd + ["shell", f"kill `cat {self._workspace}/rpc_pid.txt`"]
+        )
+
+    def _get_workspace_size(self) -> int:
+        """Get workspace base directory size in MB"""
+        line = subprocess.check_output(
+            self._adb_device_sub_cmd + ["shell", "du", "-shm", str(ANDROID_HEXAGON_TEST_BASE_DIR)],
+            encoding="utf-8",
+        )
+        return int(line.split("\t")[0])
+
+    def _workspace_remove_latest(self):
+        # Find oldest(lower number) directory
+        latest_dir = subprocess.check_output(
+            self._adb_device_sub_cmd
+            + [
+                "shell",
+                "find",
+                str(ANDROID_HEXAGON_TEST_BASE_DIR),
+                "!",
+                "-path",
+                ".",
+                "-type",
+                "d",
+                "|",
+                "sort",
+                "-n",
+                "|",
+                "head",
+                "-1",
+            ],
+            encoding="utf-8",
+        )
+        latest_dir = latest_dir.replace("\n", "").replace("\t", "")
+
+        subprocess.check_call(self._adb_device_sub_cmd + ["shell", "rm", "-rf", latest_dir])
diff --git a/python/tvm/contrib/hexagon.py b/python/tvm/contrib/hexagon/hexagon.py
similarity index 95%
rename from python/tvm/contrib/hexagon.py
rename to python/tvm/contrib/hexagon/hexagon.py
index fe256163f73cc..35136a31f3a91 100644
--- a/python/tvm/contrib/hexagon.py
+++ b/python/tvm/contrib/hexagon/hexagon.py
@@ -22,7 +22,7 @@
 import tvm
 import tvm.ir
 import tvm.contrib.cc as cc
-from .._ffi.registry import register_func
+from ..._ffi.registry import register_func
 
 
 # Linking Hexagon shared libraries.
@@ -47,16 +47,16 @@
 
 def register_linker(f):
     """Register a function that will return the path to the Hexagon linker."""
-    return register_func("tvm.contrib.hexagon.hexagon_link", f, True)
+    return register_func("tvm.contrib.hexagon.hexagon.hexagon_link", f, True)
 
 
-@register_func("tvm.contrib.hexagon.hexagon_link")
+@register_func("tvm.contrib.hexagon.hexagon.hexagon_link")
 def hexagon_link():
     """Return path to the Hexagon linker."""
     return hexagon_link_main
 
 
-@register_func("tvm.contrib.hexagon.link_shared")
+@register_func("tvm.contrib.hexagon.hexagon.link_shared")
 def link_shared(so_name, objs, **kwargs):
     """Link shared library on Hexagon using the registered Hexagon linker.
 
@@ -83,9 +83,9 @@ def to_str(s):
 
     objs = [to_str(s) for s in objs]
 
-    linker = tvm.get_global_func("tvm.contrib.hexagon.hexagon_link")()
+    linker = tvm.get_global_func("tvm.contrib.hexagon.hexagon.hexagon_link")()
     if kwargs.get("verbose"):
-        print("tvm.contrib.hexagon.link_shared:")
+        print("tvm.contrib.hexagon.hexagon.link_shared:")
         print("  Using linker:", linker)
         print("  Library name:", so_name)
         print("  Object files:", objs)
diff --git a/python/tvm/contrib/hexagon/session.py b/python/tvm/contrib/hexagon/session.py
new file mode 100644
index 0000000000000..c413c60ce5c5f
--- /dev/null
+++ b/python/tvm/contrib/hexagon/session.py
@@ -0,0 +1,75 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License.  You may obtain a copy of the License at
+#
+#   http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied.  See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+"""Defines a Session class for Hexagon devices."""
+
+from tvm import rpc as _rpc
+
+
+class Session:
+    """Hexagon Device Session
+
+    Parameters
+    ----------
+    remote_kw : dict
+        Remote configs for RPC tracker.
+
+    session_name : str
+        Hexagon RPC session name.
+    """
+
+    def __init__(
+        self,
+        remote_kw: dict,
+        session_name: str = "hexagon-rpc",
+        remote_stack_size_bytes: int = 128 * 1024,
+    ):
+        self._session_name = session_name
+        self._remote_stack_size_bytes = remote_stack_size_bytes
+        self._remote_kw = remote_kw
+        self._rpc = None
+        self.device = None
+
+    def __enter__(self):
+        if self.device:
+            # Already initialized
+            return self
+
+        tracker = _rpc.connect_tracker(self._remote_kw["host"], self._remote_kw["port"])
+        try:
+            self._rpc = tracker.request(
+                self._remote_kw["key"],
+                priority=self._remote_kw["priority"],
+                session_timeout=self._remote_kw["timeout"],
+                session_constructor_args=[
+                    "tvm.contrib.hexagon.create_hexagon_session",
+                    self._session_name,
+                    self._remote_stack_size_bytes,
+                ],
+            )
+            self.device = self._rpc.hexagon(0)
+            return self
+
+        except RuntimeError as exception:
+            raise exception
+
+    def __exit__(self, exc_type, exc_value, exc_traceback):
+        pass
+
+    def load_module(self, path: str):
+        assert isinstance(path, str), f"Invalid path type, {type(path)} != str"
+        return self._rpc.get_function("tvm.hexagon.load_module")(path)
diff --git a/src/runtime/hexagon/rpc/android/session.cc b/src/runtime/hexagon/rpc/android/session.cc
new file mode 100644
index 0000000000000..760b84ed0d8b5
--- /dev/null
+++ b/src/runtime/hexagon/rpc/android/session.cc
@@ -0,0 +1,120 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/*!
+ * \file hexagon_session.cc
+ */
+
+#include <tvm/runtime/registry.h>
+
+extern "C" {
+#include <AEEStdDef.h>
+#include <AEEStdErr.h>
+#include <HAP_farf.h>
+#include <HAP_perf.h>
+}
+
+#include <tvm/runtime/logging.h>
+
+#include <string>
+
+#include "../../../rpc/rpc_channel.h"
+#include "../../../rpc/rpc_endpoint.h"
+#include "../../../rpc/rpc_session.h"
+#include "../hexagon_rpc.h"
+
+namespace tvm {
+namespace runtime {
+namespace hexagon {
+
+class HexagonTransportChannel : public RPCChannel {
+ public:
+  explicit HexagonTransportChannel(const std::string& uri, int remote_stack_size_bytes) {
+    if (_handle != AEE_EUNKNOWN) return;
+
+    enable_unsigned_pd(true);
+    set_remote_stack_size(remote_stack_size_bytes);
+    AEEResult rc = hexagon_rpc_open(uri.c_str(), &_handle);
+    ICHECK(rc == AEE_SUCCESS) << "Hexagon RPC Open failed. URI: " << uri.c_str();
+  }
+
+  size_t Send(const void* data, size_t size) override {
+    ICHECK(_handle != AEE_EUNKNOWN) << "RPC handle is not initialized.";
+    AEEResult rc =
+        hexagon_rpc_send(_handle, static_cast<const unsigned char*>(data), static_cast<int>(size));
+    ICHECK(rc == AEE_SUCCESS) << "hexagon_rpc_send failed: " << rc;
+    return size;
+  }
+
+  size_t Recv(void* data, size_t size) override {
+    ICHECK(_handle != AEE_EUNKNOWN) << "RPC handle is not initialized.";
+    int64_t written_size = 0;
+    AEEResult rc = hexagon_rpc_receive(_handle, static_cast<unsigned char*>(data),
+                                       static_cast<int>(size), &written_size);
+    ICHECK(rc == AEE_SUCCESS) << "hexagon_rpc_receive failed: " << rc;
+    return static_cast<size_t>(written_size);
+  }
+
+  AEEResult Close() {
+    if (_handle == AEE_EUNKNOWN) return AEE_SUCCESS;
+    return hexagon_rpc_close(_handle);
+  }
+
+ private:
+  AEEResult set_remote_stack_size(int size) {
+    remote_rpc_thread_params data;
+    data.domain = CDSP_DOMAIN_ID;
+    data.prio = -1;
+    data.stack_size = size;
+    AEEResult rc = remote_session_control(FASTRPC_THREAD_PARAMS, &data, sizeof(data));
+    if (rc != AEE_SUCCESS) {
+      LOG(ERROR) << "error setting remote stack size: " << std::hex << rc << '\n';
+    }
+    return rc;
+  }
+
+  AEEResult enable_unsigned_pd(bool enable) {
+    remote_rpc_control_unsigned_module data;
+    data.domain = CDSP_DOMAIN_ID;
+    data.enable = static_cast<int>(enable);
+    AEEResult rc = remote_session_control(DSPRPC_CONTROL_UNSIGNED_MODULE, &data, sizeof(data));
+    if (rc != AEE_SUCCESS) {
+      LOG(ERROR) << "Error " << (enable ? "enabling" : "disabling") << " unsigned PD\n";
+    }
+    return rc;
+  }
+
+  remote_handle64 _handle = AEE_EUNKNOWN;
+};
+
+TVM_REGISTER_GLOBAL("tvm.contrib.hexagon.create_hexagon_session")
+    .set_body([](TVMArgs args, TVMRetValue* rv) {
+      std::string session_name = args[0];
+      int remote_stack_size_bytes = args[1];
+      HexagonTransportChannel* hexagon_channel =
+          new HexagonTransportChannel(hexagon_rpc_URI CDSP_DOMAIN, remote_stack_size_bytes);
+      std::unique_ptr<RPCChannel> channel(hexagon_channel);
+      auto ep = RPCEndpoint::Create(std::move(channel), session_name, "", NULL);
+      auto sess = CreateClientSession(ep);
+      *rv = CreateRPCSessionModule(sess);
+    });
+
+}  // namespace hexagon
+}  // namespace runtime
+}  // namespace tvm
diff --git a/src/runtime/hexagon/rpc/android_bash.sh.template b/src/runtime/hexagon/rpc/android_bash.sh.template
new file mode 100644
index 0000000000000..7bf6d773f2f78
--- /dev/null
+++ b/src/runtime/hexagon/rpc/android_bash.sh.template
@@ -0,0 +1,24 @@
+#!/bin/sh
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License.  You may obtain a copy of the License at
+#
+#   http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied.  See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+export LD_LIBRARY_PATH=.
+./tvm_rpc_android server --port=<RPC_SERVER_PORT> --tracker=<RPC_TRACKER_HOST>:<RPC_TRACKER_PORT> --key=<HEXAGON_REMOTE_DEVICE_KEY>&
+rpc_pid=$!
+
+rm -f rpc_pid.txt
+echo $rpc_pid >> rpc_pid.txt
diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc
new file mode 100644
index 0000000000000..f09223e1ad86c
--- /dev/null
+++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc
@@ -0,0 +1,271 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+extern "C" {
+#include <AEEStdDef.h>
+#include <AEEStdErr.h>
+#include <HAP_farf.h>
+#include <HAP_perf.h>
+#include <qurt_error.h>
+#include <qurt_hvx.h>
+}
+
+#include <dlfcn.h>
+#include <tvm/runtime/object.h>
+#include <tvm/runtime/packed_func.h>
+#include <tvm/runtime/registry.h>
+
+#include <algorithm>
+#include <memory>
+#include <string>
+
+#include "../../../library_module.h"
+#include "../../../minrpc/minrpc_server.h"
+#include "../../hexagon/hexagon_common.h"
+#include "hexagon_rpc.h"
+
+// TODO(mehrdadh): make this configurable.
+#define TVM_HEXAGON_RPC_BUFF_SIZE_BYTES 2 * 1024 * 1024
+
+#define TVM_LOG_CUSTOMIZE 1
+
+namespace tvm {
+namespace runtime {
+namespace hexagon {
+
+/*!
+ * \brief Hexagon IO Handler used in HexagonRPCServer(MinRPCServer).
+ *
+ * \param read_buffer The pointer to read buffer.
+ * \param read_buffer_size_bytes The read buffer size in bytes.
+ */
+class HexagonIOHandler {
+ public:
+  explicit HexagonIOHandler(uint8_t* read_buffer, size_t read_buffer_size_bytes)
+      : read_buffer_{read_buffer},
+        read_buffer_size_bytes_{read_buffer_size_bytes},
+        read_buffer_index_{0} {}
+
+  void MessageStart(size_t message_size_bytes) {}
+
+  ssize_t PosixWrite(const uint8_t* buf, size_t write_len_bytes) {
+    HEXAGON_PRINT(ALWAYS, "HexagonIOHandler PosixWrite called, write_len_bytes: %d",
+                  write_len_bytes);
+    size_t written_size = static_cast<size_t>(
+        write_buffer_.sputn(reinterpret_cast<const char*>(buf), write_len_bytes));
+    if (written_size != write_len_bytes) {
+      HEXAGON_PRINT(ALWAYS, "HexagonIOHandler written_size failed");
+    }
+    return (ssize_t)written_size;
+  }
+
+  void MessageDone() {}
+
+  ssize_t PosixRead(uint8_t* buf, size_t read_len_bytes) {
+    HEXAGON_PRINT(ALWAYS, "HexagonIOHandler PosixRead called, %d, %d", read_len_bytes,
+                  read_buffer_index_);
+
+    uint32_t bytes_to_read = 0;
+    if ((read_buffer_index_ - read_len_bytes) < 0) {
+      bytes_to_read = read_buffer_index_;
+    } else {
+      bytes_to_read = read_len_bytes;
+    }
+
+    std::memcpy(buf, read_buffer_, bytes_to_read);
+    read_buffer_ += bytes_to_read;
+    read_buffer_index_ -= bytes_to_read;
+    if (bytes_to_read != read_len_bytes) {
+      HEXAGON_PRINT(ERROR, "Error bytes_to_read (%d) < read_len_bytes (%d).", bytes_to_read,
+                    read_len_bytes);
+    }
+    return (ssize_t)bytes_to_read;
+  }
+
+  /*!
+   * \brief Set read buffer in IOHandler to data pointer.
+   * \param data The data pointer.
+   * \param data_size_bytes The size of data in bytes.
+   *
+   * \return The status
+   */
+  AEEResult SetReadBuffer(const uint8_t* data, size_t data_size_bytes) {
+    HEXAGON_PRINT(ALWAYS, "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_index_: ",
+                  data_size_bytes, read_buffer_index_);
+    if (data_size_bytes > read_buffer_size_bytes_) {
+      return AEE_EFAILED;
+    }
+    read_buffer_ = data;
+    read_buffer_index_ = data_size_bytes;
+    return AEE_SUCCESS;
+  }
+
+  /*!
+   * \brief Get pointer to the buffer that a packet has been written to.
+   * \param buf The data pointer.
+   * \param read_size_bytes The size of read in bytes.
+   *
+   * \return The size of data that is read in bytes.
+   */
+  int64_t GetWriteBuffer(uint8_t* buf, size_t read_size_bytes) {
+    HEXAGON_PRINT(ALWAYS, "HexagonIOHandler GetWriteBuffer called, read_len_bytes: %d",
+                  read_size_bytes);
+    return write_buffer_.sgetn(reinterpret_cast<char*>(buf), read_size_bytes);
+  }
+
+  void Close() { HEXAGON_PRINT(ALWAYS, "HexagonIOHandler Close called"); }
+
+  void Exit(int code) { exit(code); }
+
+ private:
+  const uint8_t* read_buffer_;
+  uint32_t read_buffer_index_;
+  size_t read_buffer_size_bytes_;
+
+  std::stringbuf write_buffer_;
+};
+
+class HexagonRPCServer {
+ public:
+  explicit HexagonRPCServer(uint8_t* receive_buffer, size_t receive_buffer_size_bytes)
+      : io_{receive_buffer, receive_buffer_size_bytes}, rpc_server_{&io_} {};
+
+  /*!
+   * \brief Wrtie to IOHandler.
+   * \param data The data pointer
+   * \param data_size_bytes The data size in bytes.
+   *
+   * \return The size of data written to IOHandler.
+   */
+  int64_t Write(const uint8_t* data, size_t data_size_bytes) {
+    if (io_.SetReadBuffer(data, data_size_bytes) != AEE_SUCCESS) {
+      return -1;
+    }
+    rpc_server_.ProcessOnePacket();
+    return (int64_t)data_size_bytes;
+  }
+
+  /*!
+   * \brief Read from IOHandler.
+   * \param buf The buffer pointer
+   * \param read_size_bytes Read request size in bytes.
+   *
+   * \return The size of data that is read in bytes.
+   */
+  int64_t Read(uint8_t* buf, size_t read_size_bytes) {
+    return io_.GetWriteBuffer(buf, read_size_bytes);
+  }
+
+ private:
+  HexagonIOHandler io_;
+  MinRPCServer<HexagonIOHandler> rpc_server_;
+};
+
+}  // namespace hexagon
+}  // namespace runtime
+}  // namespace tvm
+
+namespace {
+tvm::runtime::hexagon::HexagonRPCServer* get_hexagon_rpc_server() {
+  static tvm::runtime::hexagon::HexagonRPCServer g_hexagon_rpc_server(
+      new uint8_t[TVM_HEXAGON_RPC_BUFF_SIZE_BYTES], TVM_HEXAGON_RPC_BUFF_SIZE_BYTES);
+  return &g_hexagon_rpc_server;
+}
+}  // namespace
+
+const tvm::runtime::PackedFunc get_runtime_func(const std::string& name) {
+  if (const tvm::runtime::PackedFunc* pf = tvm::runtime::Registry::Get(name)) {
+    return *pf;
+  }
+  return tvm::runtime::PackedFunc();
+}
+
+void reset_device_api() {
+  const tvm::runtime::PackedFunc api = get_runtime_func("device_api.hexagon.v2");
+  tvm::runtime::Registry::Register("device_api.hexagon", true).set_body(api);
+}
+
+int __QAIC_HEADER(hexagon_rpc_open)(const char* uri, remote_handle64* handle) {
+  *handle = static_cast<remote_handle64>(reinterpret_cast<uintptr_t>(malloc(1)));
+  if (!*handle) {
+    HEXAGON_PRINT(ERROR, "%s: cannot allocate memory", __func__);
+    return AEE_ENOMEMORY;
+  }
+  reset_device_api();
+  get_hexagon_rpc_server();
+  return AEE_SUCCESS;
+}
+
+int __QAIC_HEADER(hexagon_rpc_close)(remote_handle64 handle) {
+  HEXAGON_PRINT(ALWAYS, "%s", __func__);
+  if (handle) {
+    free(reinterpret_cast<void*>(static_cast<uintptr_t>(handle)));
+  }
+  return AEE_SUCCESS;
+}
+
+/*!
+ * \brief Send data from Host to Hexagon over RPCSession.
+ * \param _handle The remote handle
+ * \param data The data sent to host.
+ * \param dataLen The size of the data.
+ *
+ * \return The status.
+ */
+AEEResult __QAIC_HEADER(hexagon_rpc_send)(remote_handle64 _handle, const unsigned char* data,
+                                          int dataLen) {
+  int64_t written_size = get_hexagon_rpc_server()->Write(reinterpret_cast<const uint8_t*>(data),
+                                                         static_cast<size_t>(dataLen));
+  if (written_size != dataLen) {
+    HEXAGON_PRINT(ERROR, "RPC Server Write failed, written_size (%d) != dataLen (%d)", written_size,
+                  dataLen);
+    return AEE_EFAILED;
+  }
+  return AEE_SUCCESS;
+}
+
+/*!
+ * \brief Receive data from Hexagon adn send to host over RPCSession.
+ * \param _handle The remote handle
+ * \param data The buffer for receiving data
+ * \param dataLen The size of the data that is requested to read in bytes.
+ * \param buf_written_size The size of the data that is actually read in bytes.
+ *
+ * \return The status.
+ */
+AEEResult __QAIC_HEADER(hexagon_rpc_receive)(remote_handle64 _handle, unsigned char* buf,
+                                             int bufLen, int64_t* buf_written_size) {
+  int64_t read_size =
+      get_hexagon_rpc_server()->Read(reinterpret_cast<uint8_t*>(buf), static_cast<size_t>(bufLen));
+  *buf_written_size = read_size;
+  if (read_size == bufLen) {
+    return AEE_SUCCESS;
+  } else {
+    HEXAGON_PRINT(ALWAYS, "RPC Server Read failed, read_size (%d) != dataLen (%d)", read_size,
+                  bufLen);
+    return AEE_EFAILED;
+  }
+}
+
+TVM_REGISTER_GLOBAL("tvm.hexagon.load_module")
+    .set_body([](tvm::runtime::TVMArgs args, tvm::runtime::TVMRetValue* rv) {
+      std::string soname = args[0];
+      tvm::ObjectPtr<tvm::runtime::Library> n = tvm::runtime::CreateDSOLibraryObject(soname);
+      *rv = CreateModuleFromLibrary(n, tvm::runtime::hexagon::WrapPackedFunc);
+    });
diff --git a/src/runtime/hexagon/rpc/hexagon_rpc.idl b/src/runtime/hexagon/rpc/hexagon_rpc.idl
new file mode 100644
index 0000000000000..55b8d39bcb023
--- /dev/null
+++ b/src/runtime/hexagon/rpc/hexagon_rpc.idl
@@ -0,0 +1,28 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include "remote.idl"
+#include "AEEStdDef.idl"
+
+typedef sequence<octet> buffer;
+
+interface hexagon_rpc : remote_handle64 {
+  AEEResult send(in buffer data);
+  AEEResult receive(rout buffer buf, rout int64_t buf_written_size);
+};
diff --git a/src/target/llvm/codegen_hexagon.cc b/src/target/llvm/codegen_hexagon.cc
index 0cc5c5bfe83a2..fed2ad5937ee9 100644
--- a/src/target/llvm/codegen_hexagon.cc
+++ b/src/target/llvm/codegen_hexagon.cc
@@ -836,9 +836,9 @@ runtime::Module BuildHexagon(IRModule mod, Target target) {
   std::string so_name(o_name, 0, o_name.size() - 1);
   so_name += "so";
 
-  const auto* f = tvm::runtime::Registry::Get("tvm.contrib.hexagon.link_shared");
-  ICHECK(f != nullptr) << "tvm.contrib.hexagon.link_shared does not to exist, "
-                          "do import tvm.contrib.hexagon";
+  const auto* f = tvm::runtime::Registry::Get("tvm.contrib.hexagon.hexagon.link_shared");
+  ICHECK(f != nullptr) << "tvm.contrib.hexagon.hexagon.link_shared does not to exist, "
+                          "do import tvm.contrib.hexagon.hexagon";
 
   Array<PrimExpr> o_names = {StringImm(o_name)};
   int rc = (*f)(so_name, o_names);
diff --git a/src/target/llvm/llvm_common.h b/src/target/llvm/llvm_common.h
index fcc44fb8f95c0..f31f3f698ddb5 100644
--- a/src/target/llvm/llvm_common.h
+++ b/src/target/llvm/llvm_common.h
@@ -27,7 +27,6 @@
 #ifdef _MSC_VER
 #pragma warning(disable : 4141 4291 4146 4624)
 #endif
-
 #ifdef TVM_LLVM_VERSION
 
 #include <llvm/Analysis/TargetTransformInfo.h>
@@ -72,11 +71,7 @@
 #include <llvm/Support/FileSystem.h>
 #include <llvm/Support/Host.h>
 #include <llvm/Support/MemoryBuffer.h>
-#if TVM_LLVM_VERSION >= 140
-#include <llvm/MC/TargetRegistry.h>
-#else
 #include <llvm/Support/TargetRegistry.h>
-#endif
 #include <llvm/Support/TargetSelect.h>
 #include <llvm/Support/raw_ostream.h>
 #include <llvm/Target/TargetMachine.h>
diff --git a/tests/lint/check_file_type.py b/tests/lint/check_file_type.py
index 1b45ac783c29a..964003845961a 100644
--- a/tests/lint/check_file_type.py
+++ b/tests/lint/check_file_type.py
@@ -150,6 +150,8 @@
     "apps/microtvm/reference-vm/arduino/base-box/Vagrantfile.packer-template",
     "apps/microtvm/reference-vm/zephyr/Vagrantfile",
     "apps/microtvm/reference-vm/zephyr/base-box/Vagrantfile.packer-template",
+    # Hexagon
+    "src/runtime/hexagon/rpc/android_bash.sh.template",
 }
 
 
diff --git a/tests/python/contrib/test_hexagon/conftest.py b/tests/python/contrib/test_hexagon/conftest.py
index b3bd00a082833..bd2ae7c084736 100644
--- a/tests/python/contrib/test_hexagon/conftest.py
+++ b/tests/python/contrib/test_hexagon/conftest.py
@@ -18,9 +18,114 @@
 """ Hexagon testing fixtures used to deduce testing argument
     values from testing parameters """
 
+import os
+import pytest
+
 import tvm
+from tvm import rpc
+
+HEXAGON_TOOLCHAIN = "HEXAGON_TOOLCHAIN"
+TVM_TRACKER_HOST = "TVM_TRACKER_HOST"
+TVM_TRACKER_PORT = "TVM_TRACKER_PORT"
+ANDROID_TRACKER_KEY = "ANDROID_TRACKER_KEY"
+ANDROID_REMOTE_DIR = "ANDROID_REMOTE_DIR"
 
 
 @tvm.testing.fixture
 def shape_nhwc(batch, in_channel, in_size):
     return (batch, in_size, in_size, in_channel)
+
+
+def _compose(args, decs):
+    """Helper to apply multiple markers"""
+    if len(args) > 0:
+        f = args[0]
+        for d in reversed(decs):
+            f = d(f)
+        return f
+    return decs
+
+
+def requires_hexagon_toolchain(*args):
+    _requires_hexagon_toolchain = [
+        pytest.mark.skipif(
+            os.environ.get("HEXAGON_TOOLCHAIN") == None,
+            reason="HEXAGON_TOOLCHAIN environment variable is required to run this test.",
+        ),
+    ]
+
+    return _compose(args, _requires_hexagon_toolchain)
+
+
+@tvm.testing.fixture
+def android_tracker_key():
+    return os.environ["ANDROID_TRACKER_KEY"]
+
+
+@tvm.testing.fixture
+def tvm_tracker_host():
+    return os.environ["TVM_TRACKER_HOST"]
+
+
+@tvm.testing.fixture
+def tvm_tracker_port():
+    return int(os.environ["TVM_TRACKER_PORT"])
+
+
+@tvm.testing.fixture
+def remote_path():
+    dso_binary = "test_binary.so"
+    return os.path.join(os.environ["ANDROID_REMOTE_DIR"], dso_binary)
+
+
+@tvm.testing.fixture
+def rpc_sess(android_tracker_key, tvm_tracker_host, tvm_tracker_port):
+    from tvm import rpc
+
+    tracker = rpc.connect_tracker(tvm_tracker_host, tvm_tracker_port)
+    remote = tracker.request(android_tracker_key, priority=0, session_timeout=600)
+    return remote
+
+
+def requires_rpc_tracker_and_android_key(*args):
+    """Mark a test as requiring an RPC tracker to exist in
+    the host environment to run."""
+    _requires_rpc_tracker = [
+        *tvm.testing.requires_rpc(),
+        pytest.mark.skipif(
+            os.environ.get(TVM_TRACKER_HOST) == None,
+            reason="Missing environment variable, TVM_TRACKER_HOST",
+        ),
+        pytest.mark.skipif(
+            os.environ.get(TVM_TRACKER_PORT) == None,
+            reason="Missing environment variable, TVM_TRACKER_PORT",
+        ),
+        pytest.mark.skipif(
+            os.environ.get(ANDROID_TRACKER_KEY) == None,
+            reason="Missing environment variable, ANDROID_TRACKER_KEY",
+        ),
+        pytest.mark.skipif(
+            os.environ.get(ANDROID_REMOTE_DIR) == None,
+            reason="Missing environment variable, ANDROID_REMOTE_DIR",
+        ),
+    ]
+
+    return _compose(args, _requires_rpc_tracker)
+
+
+def requires_rpc_tracker(*args):
+    """Mark a test as requiring an RPC tracker to exist in
+    the host environment to run."""
+    _requires_rpc_tracker = [
+        *tvm.testing.requires_rpc(),
+        pytest.mark.skipif(
+            os.environ.get("TVM_TRACKER_HOST") == None,
+            reason="Missing environment variable, TVM_TRACKER_HOST",
+        ),
+        pytest.mark.skipif(
+            os.environ.get("TVM_TRACKER_PORT") == None,
+            reason="Missing environment variable, TVM_TRACKER_PORT",
+        ),
+    ]
+
+    return _compose(args, _requires_rpc_tracker)
diff --git a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py b/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py
index 2b18911bacf28..839fdc9bc29d0 100644
--- a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py
+++ b/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py
@@ -16,104 +16,18 @@
 # under the License.
 
 import os
-import sys
 
 import tvm
 import tvm.testing
 from tvm import te
-import tvm.contrib.hexagon as hexagon
+import tvm.contrib.hexagon.hexagon as hexagon
 from tvm.contrib import utils
-from tvm import rpc
 import numpy as np
 
-import pytest
+from ..conftest import requires_hexagon_toolchain, requires_rpc_tracker_and_android_key
 
-HEXAGON_TOOLCHAIN = "HEXAGON_TOOLCHAIN"
-TVM_TRACKER_HOST = "TVM_TRACKER_HOST"
-TVM_TRACKER_PORT = "TVM_TRACKER_PORT"
-ANDROID_TRACKER_KEY = "ANDROID_TRACKER_KEY"
-ANDROID_REMOTE_DIR = "ANDROID_REMOTE_DIR"
 
-
-def _compose(args, decs):
-    """Helper to apply multiple markers"""
-    if len(args) > 0:
-        f = args[0]
-        for d in reversed(decs):
-            f = d(f)
-        return f
-    return decs
-
-
-def requires_hexagon_toolchain(*args):
-    _requires_rpc_tracker = [
-        *tvm.testing.requires_rpc(),
-        pytest.mark.skipif(
-            os.environ.get(HEXAGON_TOOLCHAIN) == None,
-            reason="HEXAGON_TOOLCHAIN environment variable is required to run Hexagon proxy rpc tests",
-        ),
-    ]
-
-    return _compose(args, _requires_rpc_tracker)
-
-
-def requires_rpc_tracker(*args):
-    """Mark a test as requiring an RPC tracker to exist in
-    the host environment to run."""
-    _requires_rpc_tracker = [
-        *tvm.testing.requires_rpc(),
-        pytest.mark.skipif(
-            os.environ.get(TVM_TRACKER_HOST) == None,
-            reason="Missing environment variable, TVM_TRACKER_HOST",
-        ),
-        pytest.mark.skipif(
-            os.environ.get(TVM_TRACKER_PORT) == None,
-            reason="Missing environment variable, TVM_TRACKER_PORT",
-        ),
-        pytest.mark.skipif(
-            os.environ.get(ANDROID_TRACKER_KEY) == None,
-            reason="Missing environment variable, ANDROID_TRACKER_KEY",
-        ),
-        pytest.mark.skipif(
-            os.environ.get(ANDROID_REMOTE_DIR) == None,
-            reason="Missing environment variable, ANDROID_REMOTE_DIR",
-        ),
-    ]
-
-    return _compose(args, _requires_rpc_tracker)
-
-
-@tvm.testing.fixture
-def android_tracker_key():
-    return os.environ["ANDROID_TRACKER_KEY"]
-
-
-@tvm.testing.fixture
-def tvm_tracker_host():
-    return os.environ["TVM_TRACKER_HOST"]
-
-
-@tvm.testing.fixture
-def tvm_tracker_port():
-    return int(os.environ["TVM_TRACKER_PORT"])
-
-
-@tvm.testing.fixture
-def remote_path():
-    dso_binary = "test_binary.so"
-    return os.path.join(os.environ["ANDROID_REMOTE_DIR"], dso_binary)
-
-
-@tvm.testing.fixture
-def rpc_sess(android_tracker_key, tvm_tracker_host, tvm_tracker_port):
-    from tvm import rpc
-
-    tracker = rpc.connect_tracker(tvm_tracker_host, tvm_tracker_port)
-    remote = tracker.request(android_tracker_key, priority=0, session_timeout=600)
-    return remote
-
-
-@requires_rpc_tracker
+@requires_rpc_tracker_and_android_key
 @requires_hexagon_toolchain
 class TestMatMul:
     M = tvm.testing.parameter(32)
diff --git a/tests/python/contrib/test_hexagon/rpc/__init__.py b/tests/python/contrib/test_hexagon/rpc/__init__.py
new file mode 100644
index 0000000000000..92e96bf52e8bc
--- /dev/null
+++ b/tests/python/contrib/test_hexagon/rpc/__init__.py
@@ -0,0 +1,18 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License.  You may obtain a copy of the License at
+#
+#   http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied.  See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+""" Testing infrastructure for Hexagon RPC"""
diff --git a/tests/python/contrib/test_hexagon/rpc/conftest.py b/tests/python/contrib/test_hexagon/rpc/conftest.py
new file mode 100644
index 0000000000000..50c199a2f56f0
--- /dev/null
+++ b/tests/python/contrib/test_hexagon/rpc/conftest.py
@@ -0,0 +1,31 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License.  You may obtain a copy of the License at
+#
+#   http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied.  See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+import pytest
+
+
+def pytest_addoption(parser):
+    parser.addoption(
+        "--serial-number",
+        required=True,
+        help=("Android device serial number list from 'adb' command."),
+    )
+
+
+@pytest.fixture
+def android_serial_number(request):
+    return request.config.getoption("--serial-number")
diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.md b/tests/python/contrib/test_hexagon/rpc/test_launcher.md
new file mode 100644
index 0000000000000..463b88e3f374c
--- /dev/null
+++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.md
@@ -0,0 +1,98 @@
+<!--- Licensed to the Apache Software Foundation (ASF) under one -->
+<!--- or more contributor license agreements.  See the NOTICE file -->
+<!--- distributed with this work for additional information -->
+<!--- regarding copyright ownership.  The ASF licenses this file -->
+<!--- to you under the Apache License, Version 2.0 (the -->
+<!--- "License"); you may not use this file except in compliance -->
+<!--- with the License.  You may obtain a copy of the License at -->
+
+<!---   http://www.apache.org/licenses/LICENSE-2.0 -->
+
+<!--- Unless required by applicable law or agreed to in writing, -->
+<!--- software distributed under the License is distributed on an -->
+<!--- "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY -->
+<!--- KIND, either express or implied.  See the License for the -->
+<!--- specific language governing permissions and limitations -->
+<!--- under the License. -->
+
+# HexagonLauncher
+HexagonLauncher is a class to handle interactions with an Android phone which includes Hexagon DSP to run a TVMModule(function/operation/graph) on Hexagon. HexagonLauncher reuses minRPC implementation to setup an RPC connection from host (your local machine) to Hexagon target which is passed through Android RPC server.
+
+## Build Required Tools/Libraries
+Here are the steps that are taken to prepare a runtime on a Hexagon device to test any model.
+
+- Build TVMRuntime library and C++ RPC server for Android.
+- Build minRPC server along with FastRPC for Hexagon.
+- Build TVM library with Hexagon support for host machine.
+- Build TVMRuntime library and C++ RPC server for host machine.
+
+To build these pieces, you can use a cmake command as follow.
+
+```bash
+cmake -DUSE_HEXAGON_RPC=ON \
+        -DUSE_ANDROID_TOOLCHAIN=/path/to/android-ndk/build/cmake/android.toolchain.cmake \
+        -DANDROID_PLATFORM=android-28 \
+        -DANDROID_ABI=arm64-v8a \
+        -DUSE_HEXAGON_ARCH=v65|v66|v68 \
+        -DUSE_HEXAGON_SDK=/path/to/Hexagon/SDK \
+        -DUSE_HEXAGON_TOOLCHAIN=/path/to/Hexagon/toolchain/ \
+        -DUSE_LLVM=/path/to/llvm/bin/llvm-config \
+        -DUSE_CPP_RPC=ON \
+        -DCMAKE_CXX_COMPILER=/path/to/clang++ \    
+        -DCMAKE_CXX_FLAGS='-stdlib=libc++' ..
+```
+
+## Testing Using HexagonLauncher
+Before starting a test you need to run an RPC tracker on your local machine and export HOST and PORT as environment variables. Also, you need to export Clang libraries to `LD_LIBRARY_PATH` and Hexagon toolchain to `HEXAGON_TOOLCHAIN`.
+
+```bash
+export LD_LIBRARY_PATH="$LD_LIBRARY_PATH:/path/to/clang++/lib"
+export HEXAGON_TOOLCHAIN="/path/to/Hexagon/toolchain/"
+
+export TVM_TRACKER_HOST="0.0.0.0"
+export TVM_TRACKER_PORT=9192
+python -m tvm.exec.rpc_tracker --host $TVM_TRACKER_HOST --port $TVM_TRACKER_PORT
+```
+
+Now, follow these steps to create an RPC session from host to Hexagon.
+
+```python
+# create an HexagonLauncher instance
+launcher = HexagonLauncher(serial_number="Serial number taken from `adb devices` command")
+
+# Create a workspace directory for this test on Android.
+# Upload required Android artifacts including TVMRuntime library and RPC server to Android workspace.
+# Uses port `forward` and `reverse` to open connection on certain ports that TVM uses to connect to RPC tracker.
+# Execute `android_bash.sh` on Android which creates two RPC servers and connects them to RPC tracker running on host machine. 
+launcher.android_run_rpc(rpc_tracker_host="TVM_TRACKER_HOST", rpc_tracker_port="TVM_TRACKER_PORT")
+
+# Upload Hexagon RPC libraries to Android workspace.
+launcher.hexagon_setup()
+
+# Create an RPC session from host to Hexagon.
+remote_kw = {
+    "host": "TVM_TRACKER_HOST",
+    "port": "TVM_TRACKER_PORT",
+    "priority": 0,
+    "timeout": 60,
+}
+launcher.hexagon_session_setup(remote_kw)
+
+# Upload TVMModule binary file to Android remote.
+launcher.upload("Path to DSO binary file on host", "DSO filename on Android remote")
+```
+
+- To execute a single function/operator on Hexagon, follow these steps.
+    ```python
+    # Enter session.
+    with launcher.session as sess:
+        # dlopen DSO binary file on Hexagon.
+        mod = launcher.get_module(dso_binary)
+        # Use mod to run function/operator on Hexagon...
+    ```
+- Or, follow these steps to create a GraphExecutor and run a JSON graph.
+    ```python
+    graph_mod = launcher.get_local_graph_executor(lowered, dso_binary)
+    graph_mod.set_input(...)
+    graph_mod.run(...)
+    ```
diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.py b/tests/python/contrib/test_hexagon/rpc/test_launcher.py
new file mode 100644
index 0000000000000..d705541f2b8ed
--- /dev/null
+++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.py
@@ -0,0 +1,213 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License.  You may obtain a copy of the License at
+#
+#   http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied.  See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+import sys
+import pytest
+import numpy as np
+import os
+
+import tvm.testing
+from tvm import te
+from tvm import relay
+from tvm.relay.backend import Executor, Runtime
+from tvm.contrib import utils, ndk
+from tvm.contrib.hexagon.build import HexagonLauncher
+import tvm.contrib.hexagon.hexagon as hexagon
+
+from ..conftest import requires_rpc_tracker, requires_hexagon_toolchain
+
+
+@requires_rpc_tracker
+@requires_hexagon_toolchain
+def test_add(tvm_tracker_host, tvm_tracker_port, android_serial_number):
+    dtype = "int8"
+    A = tvm.te.placeholder((2,), dtype=dtype)
+    B = tvm.te.placeholder((1,), dtype=dtype)
+    C = tvm.te.compute(A.shape, lambda i: A[i] + B[0], name="C")
+    sched = tvm.te.create_schedule(C.op)
+
+    target_hexagon = tvm.target.hexagon("v68", link_params=True)
+    func = tvm.build(
+        sched, [A, B, C], tvm.target.Target(target_hexagon, host=target_hexagon), name="add"
+    )
+
+    temp = utils.tempdir()
+    dso_binary = "test_binary.so"
+    dso_binary_path = temp.relpath(dso_binary)
+    func.save(dso_binary_path)
+
+    launcher = HexagonLauncher(serial_number=android_serial_number)
+    launcher.android_run_rpc(rpc_tracker_host=tvm_tracker_host, rpc_tracker_port=tvm_tracker_port)
+    launcher.hexagon_setup()
+    remote_kw = {
+        "host": tvm_tracker_host,
+        "port": tvm_tracker_port,
+        "priority": 0,
+        "timeout": 60,
+    }
+    launcher.hexagon_session_setup(remote_kw)
+    launcher.upload(dso_binary_path, dso_binary)
+
+    with launcher.session as sess:
+        mod = launcher.get_module(dso_binary)
+        A_data = tvm.nd.array(np.array([2, 3], dtype=dtype), device=sess.device)
+        assert (A_data.numpy() == np.array([2, 3])).all()
+        B_data = tvm.nd.array(np.array([4], dtype=dtype), device=sess.device)
+        assert (B_data.numpy() == np.array([4])).all()
+        C_data = tvm.nd.array(np.array([0, 0], dtype=dtype), device=sess.device)
+        assert (C_data.numpy() == np.array([0, 0])).all()
+
+        mod["add"](A_data, B_data, C_data)
+        assert (C_data.numpy() == np.array([6, 7])).all()
+    launcher.close()
+
+
+class TestMatMul:
+    M = tvm.testing.parameter(32)
+    N = tvm.testing.parameter(32)
+    K = tvm.testing.parameter(32)
+
+    @requires_rpc_tracker
+    @requires_hexagon_toolchain
+    def test_matmul(self, tvm_tracker_host, tvm_tracker_port, android_serial_number, M, N, K):
+        X = te.placeholder((M, K), dtype="float32")
+        Y = te.placeholder((K, N), dtype="float32")
+        k1 = te.reduce_axis((0, K), name="k1")
+        Z = te.compute((M, N), lambda i, j: te.sum(X[i, k1] * Y[k1, j], axis=[k1]))
+        schedule = te.create_schedule(Z.op)
+
+        target_hexagon = tvm.target.hexagon("v68", link_params=True)
+        func = tvm.build(
+            schedule, [X, Y, Z], tvm.target.Target(target_hexagon, host=target_hexagon)
+        )
+
+        temp = utils.tempdir()
+        dso_binary = "test_binary.so"
+        dso_binary_path = temp.relpath(dso_binary)
+        func.save(dso_binary_path)
+
+        launcher = HexagonLauncher(serial_number=android_serial_number)
+        launcher.android_run_rpc(
+            rpc_tracker_host=tvm_tracker_host, rpc_tracker_port=tvm_tracker_port
+        )
+        launcher.hexagon_setup()
+        remote_kw = {
+            "host": tvm_tracker_host,
+            "port": tvm_tracker_port,
+            "priority": 0,
+            "timeout": 60,
+        }
+        launcher.hexagon_session_setup(remote_kw)
+        launcher.upload(dso_binary_path, dso_binary)
+
+        x = np.random.uniform(size=[i.value for i in X.shape]).astype(X.dtype)
+        y = np.random.uniform(size=[i.value for i in Y.shape]).astype(Y.dtype)
+        z = np.zeros([i.value for i in Z.shape], dtype=Z.dtype)
+
+        with launcher.session as sess:
+            mod = launcher.get_module(dso_binary)
+            xt = tvm.nd.array(x, device=sess.device)
+            yt = tvm.nd.array(y, device=sess.device)
+            zt = tvm.nd.array(z, device=sess.device)
+            mod(xt, yt, zt)
+
+        target_llvm = tvm.target.Target("llvm")
+        mod = tvm.build(schedule, [X, Y, Z], tvm.target.Target(target_llvm, host=target_llvm))
+        device = tvm.cpu(0)
+        xtcpu = tvm.nd.array(x, device)
+        ytcpu = tvm.nd.array(y, device)
+        ztcpu = tvm.nd.array(z, device)
+        mod(xtcpu, ytcpu, ztcpu)
+        launcher.close()
+
+        tvm.testing.assert_allclose(zt.numpy(), ztcpu.numpy(), rtol=1e-4)
+
+
+@requires_rpc_tracker
+@requires_hexagon_toolchain
+def test_graph_executor(tvm_tracker_host, tvm_tracker_port, android_serial_number):
+    dtype = "float32"
+    data = relay.var("data", relay.TensorType((1, 64, 64, 3), dtype))
+    weight = relay.var("weight", relay.TensorType((5, 5, 3, 8), dtype))
+    y = relay.nn.conv2d(
+        data,
+        weight,
+        padding=(2, 2),
+        kernel_size=(5, 5),
+        data_layout="NHWC",
+        kernel_layout="HWIO",
+        out_dtype="float32",
+    )
+    f = relay.Function([data, weight], y)
+    relay_mod = tvm.IRModule.from_expr(f)
+    relay_mod = relay.transform.InferType()(relay_mod)
+
+    target_hexagon = tvm.target.hexagon("v68")
+    runtime = Runtime("cpp")
+    executor = Executor("graph")
+
+    temp = utils.tempdir()
+    dso_binary = "test_binary.so"
+    dso_binary_path = temp.relpath(dso_binary)
+
+    with tvm.transform.PassContext(opt_level=3):
+        lowered = tvm.relay.build(
+            relay_mod,
+            tvm.target.Target(target_hexagon, host=target_hexagon),
+            runtime=runtime,
+            executor=executor,
+        )
+        lowered.get_lib().save(dso_binary_path)
+
+    launcher = HexagonLauncher(serial_number=android_serial_number)
+    launcher.android_run_rpc(rpc_tracker_host=tvm_tracker_host, rpc_tracker_port=tvm_tracker_port)
+    launcher.hexagon_setup()
+    remote_kw = {
+        "host": tvm_tracker_host,
+        "port": tvm_tracker_port,
+        "priority": 0,
+        "timeout": 60,
+    }
+    launcher.hexagon_session_setup(remote_kw)
+    launcher.upload(dso_binary_path, dso_binary)
+
+    graph_mod = launcher.get_graph_executor(lowered, dso_binary)
+    weight_in = np.random.rand(5, 5, 3, 8).astype(dtype=dtype)
+    data_in = np.random.rand(1, 64, 64, 3).astype(dtype=dtype)
+    graph_mod.set_input(weight=weight_in)
+    graph_mod.run(data=data_in)
+    hexagon_output = graph_mod.get_output(0).numpy()
+
+    target_llvm = tvm.target.Target("llvm")
+    with tvm.transform.PassContext(opt_level=3):
+        llvm_lowered = tvm.relay.build(
+            relay_mod,
+            tvm.target.Target(target_llvm, host=target_llvm),
+            runtime=runtime,
+            executor=executor,
+        )
+    llvm_graph_mod = tvm.contrib.graph_executor.GraphModule(llvm_lowered["default"](tvm.cpu(0)))
+    llvm_graph_mod.set_input(weight=weight_in)
+    llvm_graph_mod.run(data=data_in)
+    expected_output = llvm_graph_mod.get_output(0).numpy()
+    launcher.close()
+
+    tvm.testing.assert_allclose(hexagon_output, expected_output, rtol=1e-4, atol=1e-5)
+
+
+if __name__ == "__main__":
+    sys.exit(pytest.main(sys.argv))
diff --git a/tests/python/unittest/test_target_codegen_hexagon.py b/tests/python/unittest/test_target_codegen_hexagon.py
index 220fa869d8409..ef0eb4ff5a7e8 100644
--- a/tests/python/unittest/test_target_codegen_hexagon.py
+++ b/tests/python/unittest/test_target_codegen_hexagon.py
@@ -23,12 +23,12 @@
 import tvm
 import tvm.relay
 import tvm.testing
-import tvm.contrib.hexagon as hexagon
+import tvm.contrib.hexagon.hexagon as hexagon
 
 
 @pytest.fixture(autouse=True)
 def register_linker():
-    original_linker = tvm.contrib.hexagon.hexagon_link()
+    original_linker = tvm.contrib.hexagon.hexagon.hexagon_link()
     # Register a phony linker, so that we can test codegen without a Hexagon toolchain.
     hexagon.register_linker(lambda: "/bin/true")
     yield None