diff --git a/CMakeLists.txt b/CMakeLists.txt
index 4f80d86..348728d 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -104,6 +104,7 @@
 
 add_subdirectory(samples)
 add_subdirectory(ssd_postprocess)
+add_subdirectory(vmvx_ukernel)
 
 # Add pigweed support
 include($ENV{ROOTDIR}/sw/pigweed/pw_build/pigweed.cmake)
diff --git a/README.md b/README.md
index 049ca2c..cc6bc6b 100644
--- a/README.md
+++ b/README.md
@@ -28,6 +28,7 @@
 * samples: Codegen and execution of ML models based on IREE
   * simple_vec_mul: Point-wise vector multiplication examples
 * ssd_postprocess: Vision postprocessing library for Single-Shot Detectors (SSD)
+* vmvx_ukernel: Custom microkernel library for VMVX
 
 ## Build the project
 
diff --git a/device/CMakeLists.txt b/device/CMakeLists.txt
index d7e527a..ece3cb0 100644
--- a/device/CMakeLists.txt
+++ b/device/CMakeLists.txt
@@ -25,5 +25,5 @@
     iree::hal
     iree::hal::drivers::local_sync::sync_driver
     iree::hal::local
-    iree::hal::local::loaders::vmvx_module_loader
+    vmvx_ukernel::vmvx_module_loader
 )
diff --git a/vmvx_ukernel/CMakeLists.txt b/vmvx_ukernel/CMakeLists.txt
new file mode 100644
index 0000000..8268a3e
--- /dev/null
+++ b/vmvx_ukernel/CMakeLists.txt
@@ -0,0 +1,76 @@
+set(IREE_RUNTIME_SOURCE_DIR "${IREE_SOURCE_DIR}/runtime/src/iree")
+
+iree_cc_library(
+  NAME
+    arch
+  HDRS
+    "${IREE_RUNTIME_SOURCE_DIR}/builtins/ukernel/arch/mmt4d_arch.h"
+    "${IREE_RUNTIME_SOURCE_DIR}/builtins/ukernel/arch/pack_arch.h"
+  SRCS
+    "mmt4d_arch.c"
+    "pack_arch.c"
+  DEPS
+    iree::builtins::ukernel::common
+    iree::builtins::ukernel::generic
+  PUBLIC
+)
+
+iree_cc_library(
+  NAME
+    ukernel
+  HDRS
+    "${IREE_RUNTIME_SOURCE_DIR}/builtins/ukernel/elementwise.h"
+    "${IREE_RUNTIME_SOURCE_DIR}/builtins/ukernel/mmt4d.h"
+    "${IREE_RUNTIME_SOURCE_DIR}/builtins/ukernel/pack.h"
+    "${IREE_RUNTIME_SOURCE_DIR}/builtins/ukernel/unpack.h"
+  SRCS
+    "${IREE_RUNTIME_SOURCE_DIR}/builtins/ukernel/mmt4d.c"
+    "${IREE_RUNTIME_SOURCE_DIR}/builtins/ukernel/pack.c"
+    "${IREE_RUNTIME_SOURCE_DIR}/builtins/ukernel/unpack.c"
+  DEPS
+    ::arch
+    iree::builtins::ukernel::common
+    iree::builtins::ukernel::elementwise
+  PUBLIC
+)
+
+iree_cc_library(
+  NAME
+    vmvx
+  HDRS
+    "${IREE_RUNTIME_SOURCE_DIR}/modules/vmvx/module.h"
+  TEXTUAL_HDRS
+    "${IREE_RUNTIME_SOURCE_DIR}/modules/vmvx/exports.inl"
+  SRCS
+    "${IREE_RUNTIME_SOURCE_DIR}/modules/vmvx/module.c"
+  DEFINES
+    "IREE_HAVE_VMVX_MODULE"
+  DEPS
+    ::ukernel
+    iree::base
+    iree::base::tracing
+    iree::base::internal::cpu
+    iree::vm
+  PUBLIC
+)
+
+iree_cc_library(
+  NAME
+    vmvx_module_loader
+  HDRS
+    "${IREE_RUNTIME_SOURCE_DIR}/hal/local/loaders/vmvx_module_loader.h"
+  SRCS
+    "${IREE_RUNTIME_SOURCE_DIR}/hal/local/loaders/vmvx_module_loader.c"
+  DEPS
+    ::vmvx
+    iree::base
+    iree::base::tracing
+    iree::hal
+    iree::hal::local::executable_library
+    iree::hal::local::executable_loader
+    iree::vm
+    iree::vm::bytecode_module
+  DEFINES
+    "IREE_HAVE_HAL_EXECUTABLE_LOADER_VMVX_MODULE=1"
+  PUBLIC
+)
diff --git a/vmvx_ukernel/mmt4d_arch.c b/vmvx_ukernel/mmt4d_arch.c
new file mode 100644
index 0000000..d7274ba
--- /dev/null
+++ b/vmvx_ukernel/mmt4d_arch.c
@@ -0,0 +1,25 @@
+/*
+ * Copyright 2022 Google LLC
+ *
+ * Licensed 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 "iree/builtins/ukernel/arch/mmt4d_arch.h"
+
+#include "iree/builtins/ukernel/mmt4d_generic.h"
+
+iree_uk_mmt4d_tile_func_t iree_uk_mmt4d_select_tile_func_arch(
+    const iree_uk_mmt4d_params_t* params) {
+  // TODO(lundong): to be replaced with Kelvin
+  return iree_uk_mmt4d_select_tile_func_generic(params);
+}
diff --git a/vmvx_ukernel/pack_arch.c b/vmvx_ukernel/pack_arch.c
new file mode 100644
index 0000000..b962938
--- /dev/null
+++ b/vmvx_ukernel/pack_arch.c
@@ -0,0 +1,25 @@
+/*
+ * Copyright 2022 Google LLC
+ *
+ * Licensed 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 "iree/builtins/ukernel/arch/pack_arch.h"
+
+#include "iree/builtins/ukernel/pack_generic.h"
+
+iree_uk_pack_tile_func_t iree_uk_pack_select_tile_func_arch(
+    const iree_uk_pack_params_t* params) {
+  // TODO(lundong): to be replaced with Kelvin
+  return iree_uk_pack_select_tile_func_generic(params);
+}