Repository: secretflow/spu (language: python)

Issue:
SPU项目编译错误
### Issue Type

Support

### Modules Involved

Others

### Have you reproduced the bug with SPU HEAD?

Yes

### Installation Kind

source

### SPU Version

7dcab95a83c9940b05e09dee846bb07aeff5ab94

### OS Platform and Distribution

Linux Ubuntu 22.04

### Python Version

3.10.9

### Compiler Version

11.3.0

### Current Behavior?

我运行`bazel build //... -c dbg`，遇到报错：

```
CMake Error in CMakeLists.txt:
  IMPORTED_LOCATION not set for imported target "HEXL::hexl" configuration
  "Debug".
```

完整log见附件。
[error.log](https://github.com/secretflow/spu/files/11521905/error.log)


### Standalone code to reproduce the issue

```Python
bazel build //... -c dbg
```


### Relevant log output

_No response_


Discussion / hints:

Fix patch:
diff --git a/.vscode/cspell.json b/.vscode/cspell.json
index 615a7559..074e2f89 100644
--- a/.vscode/cspell.json
+++ b/.vscode/cspell.json
@@ -55,7 +55,7 @@
         "galios",
         "galkey",
         "hashset",
-        "hctx",
+        "sctx",
         "Homomorphic",
         "iknp",
         "imgs",
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 3ee74995..4c44fb8b 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -15,6 +15,7 @@
 - [Bugfix] Fix gen cache in unbalanced psi
 - [Bugfix] Fix possible compilation cache collision
 - [Improvement] Improve accuracy of mean with large size input
+- [API] Add compiler options
 
 
 ## 20230406
diff --git a/bazel/seal.BUILD b/bazel/seal.BUILD
index 419f5481..98e972c3 100644
--- a/bazel/seal.BUILD
+++ b/bazel/seal.BUILD
@@ -21,7 +21,15 @@ filegroup(
     srcs = glob(["**"]),
 )
 
-arm64_config = {
+config_setting(
+    name = "can_use_hexl",
+    constraint_values = [
+        "@platforms//cpu:x86_64",
+    ],
+    values = {"compilation_mode": "opt"},
+)
+
+default_config = {
     "SEAL_USE_MSGSL": "OFF",
     "SEAL_BUILD_DEPS": "OFF",
     "SEAL_USE_ZSTD": "ON",
@@ -30,7 +38,7 @@ arm64_config = {
     "CMAKE_INSTALL_LIBDIR": "lib",
 }
 
-x64_config = {
+x64_hexl_config = {
     "SEAL_USE_MSGSL": "OFF",
     "SEAL_BUILD_DEPS": "OFF",
     "SEAL_USE_ZSTD": "ON",
@@ -45,8 +53,8 @@ x64_config = {
 spu_cmake_external(
     name = "seal",
     cache_entries = select({
-        "@platforms//cpu:x86_64": x64_config,
-        "//conditions:default": arm64_config,
+        ":can_use_hexl": x64_hexl_config,
+        "//conditions:default": default_config,
     }),
     lib_source = "@com_github_microsoft_seal//:all",
     out_include_dir = "include/SEAL-4.0",
diff --git a/docs/development/add_protocols.rst b/docs/development/add_protocols.rst
index 3de63308..ed4bca92 100644
--- a/docs/development/add_protocols.rst
+++ b/docs/development/add_protocols.rst
@@ -47,14 +47,14 @@ of the new protocols.
 
 Generate protocol object
 ~~~~~~~~~~~~~~~~~~~~~~~~~~~
-Each MPC protocol execution environment is abstracted as a C++ instance of an `Object <https://github.com/secretflow/spu/blob/main/libspu/mpc/object.h>`_ 
-class in SPU. SPU generates an MPC object through a factory function named `CreateCompute <https://github.com/secretflow/spu/blob/main/libspu/mpc/factory.h>`_
+Each MPC protocol execution environment is abstracted as a C++ instance of an `Object <https://github.com/secretflow/spu/blob/main/libspu/core/object.h>`_ 
+class in SPU. SPU generates an MPC object through a factory function named `CreateContext <https://github.com/secretflow/spu/blob/main/libspu/mpc/factory.h>`_
 according to the runtime config. Therefore, protocol developers need to add their functions in **CreateCompute** to generate protocol objects.
 
 .. code-block:: c++
   :caption: CreateCompute function
 
-  std::unique_ptr<Object> Factory::CreateCompute(
+  std::unique_ptr<SPUContext> Factory::CreateContext(
     const RuntimeConfig& conf,
     const std::shared_ptr<yacl::link::Context>& lctx) {
     switch (conf.protocol()) {
@@ -75,7 +75,7 @@ according to the runtime config. Therefore, protocol developers need to add thei
 
 Understand protocol object
 ~~~~~~~~~~~~~~~~~~~~~~~~~~
-SPU protocol `Object <https://github.com/secretflow/spu/blob/main/libspu/mpc/object.h>`_
+SPU protocol `Object <https://github.com/secretflow/spu/blob/main/libspu/core/object.h>`_
 may be the key concept for adding new protocols. Let's take a closer look at its design.
 The goal of Object class is to realize the generalization and flexibility of developing MPC protocols through dynamic binding. 
 An Object instance has a series of kernels and states. A kernel and a state can be regarded as a 
@@ -290,7 +290,7 @@ The second way is to write and run unittest. Some protocols do not cover all the
 or developers only want to test the functionalities of some specific MPC operations (such as addition and multiplication). 
 In these cases it is more practical to run unittest. SPU developers have construct a general test frameworks in 
 `api_test.cc <https://github.com/secretflow/spu/blob/main/libspu/mpc/api_test.cc>`_ and 
-`ab_api_test.cc <https://github.com/secretflow/spu/blob/main/libspu/mpc/common/ab_api_test.cc>`_. 
+`ab_api_test.cc <https://github.com/secretflow/spu/blob/main/libspu/mpc/ab_api_test.cc>`_. 
 Developers of new protocols need to instantiate these frameworks to test their own protocol functionalities. 
 Developers can refer to the `protocol_test.cc <https://github.com/secretflow/spu/blob/main/libspu/mpc/aby3/protocol_test.cc>`_ 
 file in the aby3 directory to learn how to write their own protocol test files.
diff --git a/examples/cpp/BUILD.bazel b/examples/cpp/BUILD.bazel
index da240c72..5eabbf99 100644
--- a/examples/cpp/BUILD.bazel
+++ b/examples/cpp/BUILD.bazel
@@ -37,7 +37,7 @@ spu_cc_library(
     srcs = ["utils.cc"],
     hdrs = ["utils.h"],
     deps = [
-        "//libspu/kernel:context",
+        "//libspu/core:context",
         "@com_google_absl//absl/strings",
         "@llvm-project//llvm:Support",
         "@yacl//yacl/link:factory",
diff --git a/examples/cpp/pir/keyword_pir_client.cc b/examples/cpp/pir/keyword_pir_client.cc
index 80ff0de2..064ce6bc 100644
--- a/examples/cpp/pir/keyword_pir_client.cc
+++ b/examples/cpp/pir/keyword_pir_client.cc
@@ -58,8 +58,8 @@ constexpr uint32_t kLinkRecvTimeout = 30 * 60 * 1000;
 int main(int argc, char **argv) {
   llvm::cl::ParseCommandLineOptions(argc, argv);
 
-  auto hctx = MakeHalContext();
-  auto link_ctx = hctx->lctx();
+  auto sctx = MakeSPUContext();
+  auto link_ctx = sctx->lctx();
 
   link_ctx->SetRecvTimeout(kLinkRecvTimeout);
 
diff --git a/examples/cpp/pir/keyword_pir_server.cc b/examples/cpp/pir/keyword_pir_server.cc
index e6960a5b..b3b2817c 100644
--- a/examples/cpp/pir/keyword_pir_server.cc
+++ b/examples/cpp/pir/keyword_pir_server.cc
@@ -60,8 +60,8 @@ int main(int argc, char **argv) {
 
   SPDLOG_INFO("setup");
 
-  auto hctx = MakeHalContext();
-  auto link_ctx = hctx->lctx();
+  auto sctx = MakeSPUContext();
+  auto link_ctx = sctx->lctx();
 
   link_ctx->SetRecvTimeout(kLinkRecvTimeout);
 
diff --git a/examples/cpp/simple_dp_psi.cc b/examples/cpp/simple_dp_psi.cc
index 52de0682..13bb0e3c 100644
--- a/examples/cpp/simple_dp_psi.cc
+++ b/examples/cpp/simple_dp_psi.cc
@@ -104,7 +104,7 @@ llvm::cl::opt<int> BucketSizeOpt("bucket_size", llvm::cl::init(1 << 20),
 int main(int argc, char** argv) {  // NOLINT
   llvm::cl::ParseCommandLineOptions(argc, argv);
 
-  auto hctx = MakeHalContext();
+  auto sctx = MakeSPUContext();
 
   auto field_list = absl::StrSplit(FieldNamesOpt.getValue(), ',');
   auto batch_provider = std::make_shared<spu::psi::CsvBatchProvider>(
@@ -123,14 +123,14 @@ int main(int argc, char** argv) {  // NOLINT
 
   std::vector<size_t> intersection_idx;
 
-  auto link_ctx = hctx->lctx();
+  auto link_ctx = sctx->lctx();
 
   link_ctx->SetThrottleWindowSize(kLinkWindowSize);
 
   link_ctx->SetRecvTimeout(kLinkRecvTimeout);
 
   if (Rank.getValue() == 0) {
-    yacl::Buffer bob_items_size_buffer = hctx->lctx()->Recv(
+    yacl::Buffer bob_items_size_buffer = sctx->lctx()->Recv(
         link_ctx->NextRank(), fmt::format("peer items number"));
     size_t bob_items_size =
         spu::psi::utils::DeserializeSize(bob_items_size_buffer);
diff --git a/examples/cpp/simple_in_memory_psi.cc b/examples/cpp/simple_in_memory_psi.cc
index 36ab4019..6bfa9fd0 100644
--- a/examples/cpp/simple_in_memory_psi.cc
+++ b/examples/cpp/simple_in_memory_psi.cc
@@ -83,13 +83,13 @@ size_t CheckPsiSize(const std::vector<std::string>& data_set) {
 int main(int argc, char** argv) {
   llvm::cl::ParseCommandLineOptions(argc, argv);
 
-  auto hctx = MakeHalContext();
+  auto sctx = MakeSPUContext();
 
   std::vector<std::string> data_set =
       CreateSampleDataset(DataSize.getValue(), Rank.getValue());
 
   std::vector<std::string> intersection =
-      spu::psi::RunEcdhPsi(hctx->lctx(), data_set, yacl::link::kAllRank);
+      spu::psi::RunEcdhPsi(sctx->lctx(), data_set, yacl::link::kAllRank);
 
   // output intersection size
   // psi result size is nearly 70% of the data_size
diff --git a/examples/cpp/simple_lr.cc b/examples/cpp/simple_lr.cc
index bf24a99f..85179127 100644
--- a/examples/cpp/simple_lr.cc
+++ b/examples/cpp/simple_lr.cc
@@ -35,7 +35,7 @@
 
 using namespace spu::kernel;
 
-spu::Value train_step(spu::HalContext* ctx, const spu::Value& x,
+spu::Value train_step(spu::SPUContext* ctx, const spu::Value& x,
                       const spu::Value& y, const spu::Value& w) {
   // Padding x
   auto padding = hal::constant(ctx, 1.0F, spu::DT_FXP, {x.shape()[0], 1});
@@ -61,7 +61,7 @@ spu::Value train_step(spu::HalContext* ctx, const spu::Value& x,
   return new_w;
 }
 
-spu::Value train(spu::HalContext* ctx, const spu::Value& x, const spu::Value& y,
+spu::Value train(spu::SPUContext* ctx, const spu::Value& x, const spu::Value& y,
                  size_t num_epoch, size_t bsize) {
   const size_t num_iter = x.shape()[0] / bsize;
   auto w = hal::constant(ctx, 0.0F, spu::DT_FXP, {x.shape()[1] + 1, 1});
@@ -87,7 +87,7 @@ spu::Value train(spu::HalContext* ctx, const spu::Value& x, const spu::Value& y,
   return w;
 }
 
-spu::Value inference(spu::HalContext* ctx, const spu::Value& x,
+spu::Value inference(spu::SPUContext* ctx, const spu::Value& x,
                      const spu::Value& weight) {
   auto padding = hal::constant(ctx, 1.0F, spu::DT_FXP, {x.shape()[0], 1});
   auto padded_x = hal::concatenate(ctx, {x, hal::seal(ctx, padding)}, 1);
@@ -124,10 +124,10 @@ llvm::cl::opt<uint32_t> BatchSize("batch_size", llvm::cl::init(21),
 llvm::cl::opt<uint32_t> NumEpoch("num_epoch", llvm::cl::init(1),
                                  llvm::cl::desc("number of epoch"));
 
-std::pair<spu::Value, spu::Value> infeed(spu::HalContext* hctx,
+std::pair<spu::Value, spu::Value> infeed(spu::SPUContext* sctx,
                                          const xt::xarray<float>& ds,
                                          bool self_has_label) {
-  spu::device::ColocatedIo cio(hctx);
+  spu::device::ColocatedIo cio(sctx);
   if (self_has_label) {
     // the last column is label.
     using namespace xt::placeholders;  // required for `_` to work
@@ -135,17 +135,17 @@ std::pair<spu::Value, spu::Value> infeed(spu::HalContext* hctx,
         xt::view(ds, xt::all(), xt::range(_, ds.shape(1) - 1));
     xt::xarray<float> dy =
         xt::view(ds, xt::all(), xt::range(ds.shape(1) - 1, _));
-    cio.hostSetVar(fmt::format("x-{}", hctx->lctx()->Rank()), dx);
+    cio.hostSetVar(fmt::format("x-{}", sctx->lctx()->Rank()), dx);
     cio.hostSetVar("label", dy);
   } else {
-    cio.hostSetVar(fmt::format("x-{}", hctx->lctx()->Rank()), ds);
+    cio.hostSetVar(fmt::format("x-{}", sctx->lctx()->Rank()), ds);
   }
   cio.sync();
 
   auto x = cio.deviceGetVar("x-0");
   // Concatnate all slices
   for (size_t idx = 1; idx < cio.getWorldSize(); ++idx) {
-    x = hal::concatenate(hctx, {x, cio.deviceGetVar(fmt::format("x-{}", idx))},
+    x = hal::concatenate(sctx, {x, cio.deviceGetVar(fmt::format("x-{}", idx))},
                          1);
   }
   auto y = cio.deviceGetVar("label");
@@ -166,19 +166,19 @@ int main(int argc, char** argv) {
     ds = xt::load_csv<float>(file, ',', SkipRows.getValue());
   }
 
-  auto hctx = MakeHalContext();
+  auto sctx = MakeSPUContext();
 
-  const auto& [x, y] = infeed(hctx.get(), ds, HasLabel.getValue());
+  const auto& [x, y] = infeed(sctx.get(), ds, HasLabel.getValue());
 
   const auto w =
-      train(hctx.get(), x, y, NumEpoch.getValue(), BatchSize.getValue());
+      train(sctx.get(), x, y, NumEpoch.getValue(), BatchSize.getValue());
 
-  const auto scores = inference(hctx.get(), x, w);
+  const auto scores = inference(sctx.get(), x, w);
 
   xt::xarray<float> revealed_labels =
-      hal::dump_public_as<float>(hctx.get(), hal::reveal(hctx.get(), y));
+      hal::dump_public_as<float>(sctx.get(), hal::reveal(sctx.get(), y));
   xt::xarray<float> revealed_scores =
-      hal::dump_public_as<float>(hctx.get(), hal::reveal(hctx.get(), scores));
+      hal::dump_public_as<float>(sctx.get(), hal::reveal(sctx.get(), scores));
 
   auto mse = MSE(revealed_labels, revealed_scores);
   std::cout << "MSE = " << mse << "\n";
diff --git a/examples/cpp/simple_pphlo.cc b/examples/cpp/simple_pphlo.cc
index ce67db19..73d8d7ed 100644
--- a/examples/cpp/simple_pphlo.cc
+++ b/examples/cpp/simple_pphlo.cc
@@ -26,7 +26,7 @@
 #include "libspu/device/pphlo/pphlo_executor.h"
 
 // This example demostrates the basic compute functionality of spu vm.
-void constant_add(spu::HalContext* hctx) {
+void constant_add(spu::SPUContext* sctx) {
   // Write the assembly, this code simple add two numbers.
   // - `%1` is a constant public integer, with dtype int32 and value 1.
   // - `%2` is a constant public integer, with dtype int32 and value 2.
@@ -44,14 +44,14 @@ func.func @main() -> () {
   // Run it, with no input and output, (since the program does not contain IO)
   spu::device::SymbolTable env;
   spu::device::pphlo::PPHloExecutor executor;
-  spu::device::execute(&executor, hctx, code, {}, {}, &env);
+  spu::device::execute(&executor, sctx, code, {}, {}, &env);
 }
 
 // This example demostrates how to pass parameters.
-void parameters(spu::HalContext* hctx) {
+void parameters(spu::SPUContext* sctx) {
   // In this example, data owner also participates the computation progress,
   // which is called "colocated mode" in spu system.
-  spu::device::ColocatedIo cio(hctx);
+  spu::device::ColocatedIo cio(sctx);
 
   if (cio.getRank() == 0) {
     // rank-0, set a float variable 3.14 as 'x' to the device.
@@ -84,18 +84,18 @@ func.func @main(%arg0: tensor<!pphlo.sec<f32>>, %arg1: tensor<!pphlo.sec<i32>>)
   // - "y" binding to the second parameter (position 1).
   // - there is no output bindings.
   spu::device::pphlo::PPHloExecutor executor;
-  spu::device::execute(&executor, hctx, code, {"x", "y"}, {},
+  spu::device::execute(&executor, sctx, code, {"x", "y"}, {},
                        &cio.deviceSymbols());
 }
 
 int main(int argc, char** argv) {
   llvm::cl::ParseCommandLineOptions(argc, argv);
 
-  auto hctx = MakeHalContext();
+  auto sctx = MakeSPUContext();
 
-  parameters(hctx.get());
+  parameters(sctx.get());
 
-  constant_add(hctx.get());
+  constant_add(sctx.get());
 
   return 0;
 }
diff --git a/examples/cpp/simple_psi.cc b/examples/cpp/simple_psi.cc
index 73156f03..1ab8f687 100644
--- a/examples/cpp/simple_psi.cc
+++ b/examples/cpp/simple_psi.cc
@@ -57,7 +57,7 @@ llvm::cl::opt<double> DPPsiEpsilonOpt("epsilon", llvm::cl::init(3),
 int main(int argc, char** argv) {
   llvm::cl::ParseCommandLineOptions(argc, argv);
 
-  auto hctx = MakeHalContext();
+  auto sctx = MakeSPUContext();
 
   auto field_list = absl::StrSplit(FieldNamesOpt.getValue(), ',');
 
@@ -83,11 +83,11 @@ int main(int argc, char** argv) {
   config.set_curve_type(spu::psi::CurveType::CURVE_25519);
 
   try {
-    spu::psi::BucketPsi bucket_psi(config, hctx->lctx());
+    spu::psi::BucketPsi bucket_psi(config, sctx->lctx());
     auto report = bucket_psi.Run();
 
     SPDLOG_INFO("rank:{} original_count:{} intersection_count:{}",
-                hctx->lctx()->Rank(), report.original_count(),
+                sctx->lctx()->Rank(), report.original_count(),
                 report.intersection_count());
   } catch (const std::exception& e) {
     SPDLOG_ERROR("run psi failed: {}", e.what());
diff --git a/examples/cpp/utils.cc b/examples/cpp/utils.cc
index 92b6aaf4..829301f0 100644
--- a/examples/cpp/utils.cc
+++ b/examples/cpp/utils.cc
@@ -46,7 +46,7 @@ std::shared_ptr<yacl::link::Context> MakeLink(const std::string& parties,
   return lctx;
 }
 
-std::unique_ptr<spu::HalContext> MakeHalContext() {
+std::unique_ptr<spu::SPUContext> MakeSPUContext() {
   auto lctx = MakeLink(Parties.getValue(), Rank.getValue());
 
   spu::RuntimeConfig config;
@@ -55,5 +55,5 @@ std::unique_ptr<spu::HalContext> MakeHalContext() {
 
   config.set_enable_action_trace(EngineTrace.getValue());
   config.set_enable_type_checker(EngineTrace.getValue());
-  return std::make_unique<spu::HalContext>(config, lctx);
+  return std::make_unique<spu::SPUContext>(config, lctx);
 }
diff --git a/examples/cpp/utils.h b/examples/cpp/utils.h
index bcfba67b..bbe7ad2c 100644
--- a/examples/cpp/utils.h
+++ b/examples/cpp/utils.h
@@ -16,7 +16,7 @@
 
 #include "llvm/Support/CommandLine.h"
 
-#include "libspu/kernel/context.h"
+#include "libspu/core/context.h"
 
 extern llvm::cl::opt<std::string> Parties;
 extern llvm::cl::opt<uint32_t> Rank;
@@ -29,4 +29,4 @@ std::shared_ptr<yacl::link::Context> MakeLink(const std::string& parties,
                                               size_t rank);
 
 // Create an evaluator with setting flags.
-std::unique_ptr<spu::HalContext> MakeHalContext();
+std::unique_ptr<spu::SPUContext> MakeSPUContext();
diff --git a/examples/python/ml/stax_nn/models.py b/examples/python/ml/stax_nn/models.py
index 27d9eb85..b56fd15a 100644
--- a/examples/python/ml/stax_nn/models.py
+++ b/examples/python/ml/stax_nn/models.py
@@ -47,11 +47,11 @@ def secureml():
 def minionn():
     nn_init, nn_apply = stax.serial(
         Conv(out_chan=16, filter_shape=(5, 5), strides=(1, 1), padding='valid'),
-        Relu,
         MaxPool(window_shape=(2, 2), strides=(2, 2)),
-        Conv(out_chan=16, filter_shape=(5, 5), strides=(1, 1), padding='valid'),
         Relu,
+        Conv(out_chan=16, filter_shape=(5, 5), strides=(1, 1), padding='valid'),
         MaxPool(window_shape=(2, 2), strides=(2, 2)),
+        Relu,
         Flatten,
         Dense(100),
         Relu,
@@ -65,11 +65,11 @@ def minionn():
 def lenet():
     nn_init, nn_apply = stax.serial(
         Conv(out_chan=20, filter_shape=(5, 5), strides=(1, 1), padding='valid'),
-        Relu,
         MaxPool(window_shape=(2, 2), strides=(2, 2)),
-        Conv(out_chan=50, filter_shape=(5, 5), strides=(1, 1), padding='valid'),
         Relu,
+        Conv(out_chan=50, filter_shape=(5, 5), strides=(1, 1), padding='valid'),
         MaxPool(window_shape=(2, 2), strides=(2, 2)),
+        Relu,
         Flatten,
         Dense(500),
         Relu,
diff --git a/libspu/compiler/common/BUILD.bazel b/libspu/compiler/common/BUILD.bazel
index 71056d85..eef7efaf 100644
--- a/libspu/compiler/common/BUILD.bazel
+++ b/libspu/compiler/common/BUILD.bazel
@@ -33,6 +33,7 @@ spu_cc_library(
     hdrs = ["compilation_context.h"],
     deps = [
         ":ir_printer_config",
+        "//libspu:spu_cc_proto",
         "//libspu/core:prelude",
     ],
 )
diff --git a/libspu/compiler/common/compilation_context.cc b/libspu/compiler/common/compilation_context.cc
index 3d25a9a5..4819824c 100644
--- a/libspu/compiler/common/compilation_context.cc
+++ b/libspu/compiler/common/compilation_context.cc
@@ -49,10 +49,6 @@ CompilationContext::getIRPrinterConfig() const {
   return std::make_unique<mlir::pphlo::IRPrinterConfig>(*config);
 }
 
-void CompilationContext::enablePrettyPrintWithDir(std::string_view dir) {
-  pp_config_ = std::make_unique<mlir::pphlo::IRPrinterConfig>(dir);
-}
-
 void CompilationContext::setupPrettyPrintConfigurations(mlir::PassManager *pm) {
   if (hasPrettyPrintEnabled()) {
     getMLIRContext()->disableMultithreading();
@@ -66,4 +62,16 @@ std::filesystem::path CompilationContext::getPrettyPrintDir() const {
       ->GetPrettyPrintDir();
 }
 
+void CompilationContext::setCompilerOptions(
+    const std::string &serialized_copts) {
+  auto status = options_.ParseFromString(serialized_copts);
+
+  SPU_ENFORCE(status, "Parse compiler options failed");
+
+  if (options_.enable_pretty_print()) {
+    pp_config_ = std::make_unique<mlir::pphlo::IRPrinterConfig>(
+        options_.pretty_print_dump_dir());
+  }
+}
+
 } // namespace spu::compiler
diff --git a/libspu/compiler/common/compilation_context.h b/libspu/compiler/common/compilation_context.h
index 147c08ee..de039d21 100644
--- a/libspu/compiler/common/compilation_context.h
+++ b/libspu/compiler/common/compilation_context.h
@@ -22,6 +22,8 @@
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/Pass/PassManager.h"
 
+#include "libspu/spu.pb.h"
+
 namespace mlir {
 class PassManager;
 }
@@ -36,18 +38,14 @@ class CompilationContext {
 
   mlir::MLIRContext *getMLIRContext() { return &context_; }
 
-  /// Enable pretty print and set folder
-  /// If dir does not exist, this api will create new folder
-  void enablePrettyPrintWithDir(std::string_view dir);
-
   /// Setup pretty print for a pass manager
   void setupPrettyPrintConfigurations(mlir::PassManager *pm);
 
-  void setInputVisibilityString(const std::string &vis) { input_vis_ = vis; }
+  void setCompilerOptions(const std::string &serialized_copts);
 
-  std::string getInputVisibilityString() const { return input_vis_; }
+  const CompilerOptions &getCompilerOptions() const { return options_; }
 
-  bool hasPrettyPrintEnabled() const { return pp_config_ != nullptr; }
+  bool hasPrettyPrintEnabled() const { return options_.enable_pretty_print(); }
 
   std::filesystem::path getPrettyPrintDir() const;
 
@@ -58,7 +56,7 @@ class CompilationContext {
   mlir::MLIRContext context_;
   std::unique_ptr<mlir::PassManager::IRPrinterConfig> pp_config_;
 
-  std::string input_vis_;
+  CompilerOptions options_;
 };
 
 } // namespace spu::compiler
diff --git a/libspu/compiler/compile.cc b/libspu/compiler/compile.cc
index 57811e65..0d93b1b0 100644
--- a/libspu/compiler/compile.cc
+++ b/libspu/compiler/compile.cc
@@ -29,11 +29,11 @@
 
 namespace spu::compiler {
 
-std::string compile(CompilationContext *ctx, const std::string &serialized_ir,
-                    const std::string &ir_type) {
+std::string compile(CompilationContext *ctx,
+                    const std::string &serialized_source) {
   // Call front end
   FE fe(ctx);
-  auto mlir_module = fe.doit(serialized_ir, ir_type);
+  auto mlir_module = fe.doit(serialized_source);
 
   // Run core passes
   Core core(ctx);
diff --git a/libspu/compiler/compile.h b/libspu/compiler/compile.h
index c4ae5e60..ae1d36b3 100644
--- a/libspu/compiler/compile.h
+++ b/libspu/compiler/compile.h
@@ -14,16 +14,13 @@
 
 #pragma once
 
-#include <filesystem>
 #include <string>
-#include <utility>
-#include <vector>
 
 namespace spu::compiler {
 
 class CompilationContext;
 
-std::string compile(CompilationContext *ctx, const std::string &serialized_ir,
-                    const std::string &ir_type);
+std::string compile(CompilationContext *ctx,
+                    const std::string &serialized_source);
 
 } // namespace spu::compiler
diff --git a/libspu/compiler/core/core.cc b/libspu/compiler/core/core.cc
index 43338e62..7651c09d 100644
--- a/libspu/compiler/core/core.cc
+++ b/libspu/compiler/core/core.cc
@@ -42,23 +42,40 @@ void Core::doit(mlir::ModuleOp module) {
 }
 
 void Core::buildPipeline(mlir::PassManager *pm) {
+  const auto &options = ctx_->getCompilerOptions();
+
   // lowering
   auto &optPM = pm->nest<mlir::func::FuncOp>();
-  optPM.addPass(mlir::pphlo::createOptimizeMaxPoolingPass());
+  if (!options.disable_maxpooling_optimization()) {
+    optPM.addPass(mlir::pphlo::createOptimizeMaxPoolingPass());
+  }
   optPM.addPass(mlir::pphlo::createDecomposeComparisonPass());
   optPM.addPass(mlir::pphlo::createDecomposeMinMaxPass());
-  optPM.addPass(mlir::pphlo::createOptimizeSqrtPlusEps());
-  optPM.addPass(mlir::pphlo::createRewriteDivSqrtPatterns());
+
+  if (!options.disable_sqrt_plus_epsilon_rewrite()) {
+    optPM.addPass(mlir::pphlo::createOptimizeSqrtPlusEps());
+  }
+  if (!options.disable_div_sqrt_rewrite()) {
+    optPM.addPass(mlir::pphlo::createRewriteDivSqrtPatterns());
+  }
+
   optPM.addPass(mlir::pphlo::createExpandSecretGatherPass());
 
   optPM.addPass(mlir::createCSEPass());
 
-  optPM.addPass(mlir::pphlo::createReduceTruncationPass());
-  optPM.addPass(mlir::pphlo::createLowerMixedTypeOpPass());
+  if (!options.disable_reduce_truncation_optimization()) {
+    optPM.addPass(mlir::pphlo::createReduceTruncationPass());
+  }
+
+  if (!options.disallow_mix_types_opts()) {
+    optPM.addPass(mlir::pphlo::createLowerMixedTypeOpPass());
+  }
 
   optPM.addPass(mlir::createCanonicalizerPass());
 
-  optPM.addPass(mlir::pphlo::createOptimizeSelectPass());
+  if (!options.disable_select_optimization()) {
+    optPM.addPass(mlir::pphlo::createOptimizeSelectPass());
+  }
 
   optPM.addPass(mlir::createLoopInvariantCodeMotionPass());
   optPM.addPass(mlir::createCSEPass());
diff --git a/libspu/compiler/core/core.h b/libspu/compiler/core/core.h
index e903f0f9..78dfde02 100644
--- a/libspu/compiler/core/core.h
+++ b/libspu/compiler/core/core.h
@@ -34,7 +34,7 @@ class Core final {
 private:
   CompilationContext *ctx_;
 
-  static void buildPipeline(mlir::PassManager *pm);
+  void buildPipeline(mlir::PassManager *pm);
 };
 
 } // namespace spu::compiler
diff --git a/libspu/compiler/front_end/fe.cc b/libspu/compiler/front_end/fe.cc
index f8c879ba..54fabd99 100644
--- a/libspu/compiler/front_end/fe.cc
+++ b/libspu/compiler/front_end/fe.cc
@@ -39,22 +39,39 @@ FE::FE(CompilationContext *ctx) : ctx_(ctx) {
                     mlir::func::FuncDialect>();
 }
 
-mlir::OwningOpRef<mlir::ModuleOp> FE::doit(const std::string &input,
-                                           const std::string &type) {
+mlir::OwningOpRef<mlir::ModuleOp> FE::doit(const std::string &source_str) {
+  CompilationSource source;
+  source.ParseFromString(source_str);
+
   mlir::OwningOpRef<mlir::ModuleOp> module;
-  if (type == "hlo") { // Import hlo
+  switch (source.ir_type()) {
+  case spu::SourceIRType::XLA: {
     HloImporter importer(ctx_);
-    module = importer.parseXlaModuleFromString(input);
-  } else if (type == "mhlo") { // Import mhlo
-    module =
-        mlir::parseSourceString<mlir::ModuleOp>(input, ctx_->getMLIRContext());
-  } else {
-    SPU_THROW("Unsupported input IR type");
+    module = importer.parseXlaModuleFromString(source.ir_txt());
+    break;
+  }
+  case spu::SourceIRType::MLIR_HLO: {
+    module = mlir::parseSourceString<mlir::ModuleOp>(source.ir_txt(),
+                                                     ctx_->getMLIRContext());
+    break;
+  }
+  default: {
+    SPU_THROW("Unsupported input IR type = {}", source.ir_type());
+  }
+  }
+
+  std::string input_vis_str;
+  {
+    std::vector<std::string> input_vis;
+    for (const auto &v : source.input_visibility()) {
+      input_vis.emplace_back(Visibility_Name(v));
+    }
+    input_vis_str = fmt::format("input_vis_list={}", fmt::join(input_vis, ","));
   }
 
   // Run pipeline
   mlir::PassManager pm(ctx_->getMLIRContext());
-  buildFrontEndPipeline(&pm);
+  buildFrontEndPipeline(&pm, input_vis_str);
 
   ctx_->setupPrettyPrintConfigurations(&pm);
 
@@ -67,7 +84,7 @@ mlir::OwningOpRef<mlir::ModuleOp> FE::doit(const std::string &input,
   return module;
 }
 
-void FE::buildFrontEndPipeline(mlir::PassManager *pm) {
+void FE::buildFrontEndPipeline(mlir::PassManager *pm, const std::string &args) {
 
   // mhlo side
   {
@@ -89,11 +106,12 @@ void FE::buildFrontEndPipeline(mlir::PassManager *pm) {
 
   // stablehlo now
   // Dialect conversion
-  auto vis_str = ctx_->getInputVisibilityString();
-  if (vis_str.empty()) {
-    pm->addPass(mlir::pphlo::createLegalizeToPPHloPass());
-  } else {
-    pm->addPass(mlir::pphlo::createLegalizeToPPHloPass(vis_str));
+  {
+    auto l = mlir::pphlo::createLegalizeToPPHloPass();
+    if (!args.empty()) {
+      SPU_ENFORCE(l->initializeOptions(args).succeeded());
+    }
+    pm->addPass(std::move(l));
   }
   auto &optPM = pm->nest<mlir::func::FuncOp>();
   optPM.addPass(mlir::pphlo::createLowerConversionCastPass());
diff --git a/libspu/compiler/front_end/fe.h b/libspu/compiler/front_end/fe.h
index 95efd845..0928c8a4 100644
--- a/libspu/compiler/front_end/fe.h
+++ b/libspu/compiler/front_end/fe.h
@@ -33,11 +33,10 @@ class FE final {
 public:
   explicit FE(CompilationContext *ctx);
 
-  mlir::OwningOpRef<mlir::ModuleOp> doit(const std::string &input,
-                                         const std::string &type);
+  mlir::OwningOpRef<mlir::ModuleOp> doit(const std::string &source);
 
 private:
-  void buildFrontEndPipeline(mlir::PassManager *pm);
+  void buildFrontEndPipeline(mlir::PassManager *pm, const std::string &args);
 
   CompilationContext *ctx_;
 };
diff --git a/libspu/compiler/passes/BUILD.bazel b/libspu/compiler/passes/BUILD.bazel
index 9c5e66ce..a88347d3 100644
--- a/libspu/compiler/passes/BUILD.bazel
+++ b/libspu/compiler/passes/BUILD.bazel
@@ -94,6 +94,7 @@ spu_cc_library(
         ":map_stablehlo_to_pphlo_op",
         ":pass_details",
         ":visibility_inference",
+        "//libspu/compiler/common:compilation_context",
         "//libspu/core:prelude",
         "//libspu/dialect:pphlo_dialect",
         "@llvm-project//mlir:IR",
diff --git a/libspu/compiler/passes/hlo_legalize_to_pphlo.cc b/libspu/compiler/passes/hlo_legalize_to_pphlo.cc
index 2b3577ef..d5288933 100644
--- a/libspu/compiler/passes/hlo_legalize_to_pphlo.cc
+++ b/libspu/compiler/passes/hlo_legalize_to_pphlo.cc
@@ -21,9 +21,6 @@
 #include <utility>
 #include <vector>
 
-#include "llvm/Support/JSON.h"
-#include "llvm/Support/Path.h"
-#include "llvm/Support/raw_os_ostream.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/PatternMatch.h"
@@ -31,6 +28,7 @@
 #include "mlir/Transforms/DialectConversion.h"
 #include "stablehlo/dialect/StablehloOps.h"
 
+#include "libspu/compiler/common/compilation_context.h"
 #include "libspu/compiler/passes/map_stablehlo_to_pphlo_op.h"
 #include "libspu/compiler/passes/pass_details.h"
 #include "libspu/compiler/passes/value_visibility_map.h"
@@ -44,38 +42,23 @@
 namespace mlir::pphlo {
 namespace {
 
-/// This struct carries information of io visibility
-struct IoVisibilityInfo {
-  std::vector<Visibility> inputs;
-
-  void convertFromStrings(llvm::ArrayRef<std::string> data) {
-    for (const auto &s : data) {
-      const auto symbolized = symbolizeEnum<Visibility>(s);
-      SPU_ENFORCE(symbolized.has_value());
-      inputs.emplace_back(*symbolized);
-    }
-  }
-
-  Visibility getInputVisibility(size_t idx) const {
-    if (idx >= inputs.size()) {
-      return Visibility::VIS_PUBLIC;
-    }
-    return inputs[idx];
-  }
-};
-
-ValueVisibilityMap VisibilityDiscovery(ModuleOp op,
-                                       const IoVisibilityInfo &input_vis) {
+ValueVisibilityMap
+VisibilityDiscovery(const llvm::ArrayRef<std::string> input_vis_list,
+                    ModuleOp op) {
   // Get the main function
   auto entry_func = op.lookupSymbol<mlir::func::FuncOp>("main");
 
-  SPU_ENFORCE(entry_func != nullptr);
+  SPU_ENFORCE(entry_func != nullptr, "Cannot find main entry point");
 
   ValueVisibilityMap vis_map;
   // Populate top level io visibility
   for (const auto &blockargs : entry_func.getBody().getArguments()) {
-    vis_map.setValueVisibility(
-        blockargs, input_vis.getInputVisibility(blockargs.getArgNumber()));
+    SPU_ENFORCE(blockargs.getArgNumber() < input_vis_list.size(),
+                "Input visibility list does not match actual inputs.");
+    auto v =
+        symbolizeEnum<Visibility>(input_vis_list[blockargs.getArgNumber()]);
+    SPU_ENFORCE(v.has_value(), "Input visibility list has invalid value.");
+    vis_map.setValueVisibility(blockargs, *v);
   }
 
   VisibilityInference inference(vis_map);
@@ -1407,42 +1390,11 @@ struct HloLegalizeToPPHlo
         HloToPPHloOpConverter<stablehlo::XorOp>>(converter, context, vis_map);
   }
 
-  IoVisibilityInfo vis_info_;
-
-  void parseVisibilityString() {
-    if (io_visibility_json_.empty()) {
-      return;
-    }
-    llvm::raw_os_ostream os(std::cout);
-    if (auto json_v = llvm::json::parse(io_visibility_json_)) {
-      llvm::json::Path::Root r;
-      llvm::json::ObjectMapper map(*json_v, r);
-      std::vector<std::string> str_vis;
-      if (map && map.map("inputs", str_vis)) {
-        vis_info_.convertFromStrings(str_vis);
-      } else {
-        r.printErrorContext(*json_v, os);
-      }
-    } else {
-      handleAllErrors(json_v.takeError(), [&](const llvm::ErrorInfoBase &e) {
-        os << "Failed to parse visibility JSON >>> " << io_visibility_json_
-           << " <<<: " << e.message();
-      });
-    }
-    os.flush();
-  }
-
 public:
   HloLegalizeToPPHlo(const HloLegalizeToPPHlo &) = default;
   HloLegalizeToPPHlo() = default;
-  explicit HloLegalizeToPPHlo(const std::string &io_visibility_json) {
-    io_visibility_json_ = io_visibility_json;
-  }
 
   void runOnOperation() override {
-    // This is a must step for cli workflow
-    parseVisibilityString();
-
     auto &context = getContext();
 
     RewritePatternSet patterns(&context);
@@ -1469,7 +1421,8 @@ struct HloLegalizeToPPHlo
         });
 
     // Stage 1: Run a visibility discover pass to tag all Values' visibility
-    ValueVisibilityMap vis_map = VisibilityDiscovery(getOperation(), vis_info_);
+    ValueVisibilityMap vis_map =
+        VisibilityDiscovery(input_vis_list_, getOperation());
 
     // Stage 2: Do an actual dialect conversion.
     populateHLOToPPHloConversionPattern(converter, patterns, vis_map);
@@ -1482,11 +1435,6 @@ struct HloLegalizeToPPHlo
 };
 } // namespace
 
-std::unique_ptr<OperationPass<ModuleOp>>
-createLegalizeToPPHloPass(const std::string &io_visibility_json) {
-  return std::make_unique<HloLegalizeToPPHlo>(io_visibility_json);
-}
-
 std::unique_ptr<OperationPass<ModuleOp>> createLegalizeToPPHloPass() {
   return std::make_unique<HloLegalizeToPPHlo>();
 }
diff --git a/libspu/compiler/passes/pass_details.h b/libspu/compiler/passes/pass_details.h
index 6dccf126..3960ebf8 100644
--- a/libspu/compiler/passes/pass_details.h
+++ b/libspu/compiler/passes/pass_details.h
@@ -17,6 +17,7 @@
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 
 #include "libspu/compiler/passes/passes.h"
+#include "libspu/dialect/pphlo_base_enums.h"
 #include "libspu/dialect/pphlo_dialect.h"
 
 namespace mlir::pphlo {
diff --git a/libspu/compiler/passes/passes.h b/libspu/compiler/passes/passes.h
index 65bcbc91..40be508e 100644
--- a/libspu/compiler/passes/passes.h
+++ b/libspu/compiler/passes/passes.h
@@ -16,8 +16,11 @@
 
 #include <memory>
 
-namespace mlir {
+namespace spu::compiler {
+class CompilationContext;
+}
 
+namespace mlir {
 namespace func {
 class FuncOp;
 }
@@ -29,13 +32,9 @@ class OperationPass;
 
 namespace pphlo {
 
-/// Lowers from HLO dialect to pphlo dialect with cli io_vis
+/// Lowers from HLO dialect to pphlo dialect
 std::unique_ptr<OperationPass<ModuleOp>> createLegalizeToPPHloPass();
 
-/// Lower from HLO dialect to pphlo dialect with a json io_vis string
-std::unique_ptr<OperationPass<ModuleOp>>
-createLegalizeToPPHloPass(const std::string &io_visibility_json);
-
 // Decompose comparison into lower ops when possible
 std::unique_ptr<OperationPass<func::FuncOp>> createDecomposeComparisonPass();
 
diff --git a/libspu/compiler/passes/passes.td b/libspu/compiler/passes/passes.td
index cb1cca3f..9b9353a3 100644
--- a/libspu/compiler/passes/passes.td
+++ b/libspu/compiler/passes/passes.td
@@ -23,8 +23,7 @@ def HloLegalizeToPPHloPass : Pass<"hlo-legalize-to-pphlo", "ModuleOp"> {
   let constructor = "createLegalizeToPPHloPass()";
   let dependentDialects = ["pphlo::PPHloDialect"];
   let options = [
-    Option<"io_visibility_json_", "io-visibility-json", "std::string", "",
-           "io visibility of this module">,
+    ListOption<"input_vis_list_", "input_vis_list", "std::string", "input visibilities to entry point function">,
   ];
 }
 
diff --git a/libspu/compiler/tools/mlir-pphlo-opt.cc b/libspu/compiler/tools/mlir-pphlo-opt.cc
index 07460637..0201acb8 100644
--- a/libspu/compiler/tools/mlir-pphlo-opt.cc
+++ b/libspu/compiler/tools/mlir-pphlo-opt.cc
@@ -32,6 +32,6 @@ int main(int argc, char **argv) {
   registry.insert<mlir::mhlo::MhloDialect, mlir::stablehlo::StablehloDialect,
                   mlir::pphlo::PPHloDialect, mlir::func::FuncDialect>();
 
-  return static_cast<int>(failed(
-      mlir::MlirOptMain(argc, argv, "MLIR pphlo pass driver\n", registry)));
+  return mlir::asMainReturnCode(
+      mlir::MlirOptMain(argc, argv, "MLIR pphlo pass driver\n", registry));
 }
diff --git a/libspu/core/BUILD.bazel b/libspu/core/BUILD.bazel
index e880a7df..17113531 100644
--- a/libspu/core/BUILD.bazel
+++ b/libspu/core/BUILD.bazel
@@ -321,3 +321,62 @@ spu_cc_binary(
         "@com_github_google_benchmark//:benchmark_main",
     ],
 )
+
+spu_cc_library(
+    name = "object",
+    srcs = ["object.cc"],
+    hdrs = ["object.h"],
+    deps = [
+        "//libspu/core:cexpr",
+        "//libspu/core:type",
+        "//libspu/core:value",
+    ],
+)
+
+spu_cc_library(
+    name = "cexpr",
+    srcs = ["cexpr.cc"],
+    hdrs = ["cexpr.h"],
+    deps = [
+        "//libspu/core:prelude",
+    ],
+)
+
+spu_cc_test(
+    name = "cexpr_test",
+    srcs = ["cexpr_test.cc"],
+    deps = [
+        ":cexpr",
+    ],
+)
+
+spu_cc_library(
+    name = "context",
+    srcs = ["context.cc"],
+    hdrs = ["context.h"],
+    deps = [
+        "//libspu/core:config",
+        "//libspu/core:object",
+        "//libspu/core:trace",
+        "@yacl//yacl/link",
+    ],
+)
+
+spu_cc_library(
+    name = "value",
+    srcs = ["value.cc"],
+    hdrs = ["value.h"],
+    deps = [
+        "//libspu/core:ndarray_ref",
+        "//libspu/core:prelude",
+        "//libspu/core:shape_util",
+    ],
+)
+
+spu_cc_test(
+    name = "value_test",
+    srcs = ["value_test.cc"],
+    deps = [
+        "//libspu/core:value",
+    ],
+)
diff --git a/libspu/mpc/utils/cexpr.cc b/libspu/core/cexpr.cc
similarity index 97%
rename from libspu/mpc/utils/cexpr.cc
rename to libspu/core/cexpr.cc
index bafd6755..4d715952 100644
--- a/libspu/mpc/utils/cexpr.cc
+++ b/libspu/core/cexpr.cc
@@ -12,7 +12,7 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "libspu/mpc/utils/cexpr.h"
+#include "libspu/core/cexpr.h"
 
 #include <cmath>
 #include <functional>
@@ -23,7 +23,7 @@
 
 #include "libspu/core/prelude.h"
 
-namespace spu::mpc::cexpr {
+namespace spu::ce {
 namespace {
 
 class ConstantExpr : public BaseExpr {
@@ -145,4 +145,4 @@ CExpr operator*(const CExpr& x, const CExpr& y) { return Mul(x, y); }
 CExpr operator*(const CExpr& x, Value y) { return Mul(x, Const(y)); }
 CExpr operator*(Value x, const CExpr& y) { return Mul(Const(x), y); }
 
-}  // namespace spu::mpc::cexpr
+}  // namespace spu::ce
diff --git a/libspu/mpc/utils/cexpr.h b/libspu/core/cexpr.h
similarity index 94%
rename from libspu/mpc/utils/cexpr.h
rename to libspu/core/cexpr.h
index 019a6fb5..eaa31849 100644
--- a/libspu/mpc/utils/cexpr.h
+++ b/libspu/core/cexpr.h
@@ -19,7 +19,7 @@
 #include <string>
 
 // The complexity expression library
-namespace spu::mpc::cexpr {
+namespace spu::ce {
 
 using Value = size_t;
 using Params = std::map<std::string, Value>;
@@ -59,8 +59,4 @@ CExpr operator*(const CExpr& x, const CExpr& y);
 CExpr operator*(const CExpr& x, Value y);
 CExpr operator*(Value x, const CExpr& y);
 
-}  // namespace spu::mpc::cexpr
-
-namespace spu {
-namespace ce = mpc::cexpr;
-}
+}  // namespace spu::ce
diff --git a/libspu/core/config.cc b/libspu/core/config.cc
index f792d075..e0559ca3 100644
--- a/libspu/core/config.cc
+++ b/libspu/core/config.cc
@@ -79,7 +79,7 @@ void populateRuntimecConfig(RuntimeConfig& cfg) {
   }
 
   if (cfg.sigmoid_mode() == RuntimeConfig::SIGMOID_DEFAULT) {
-    cfg.set_sigmoid_mode(RuntimeConfig::SIGMOID_MM1);
+    cfg.set_sigmoid_mode(RuntimeConfig::SIGMOID_REAL);
   }
 }
 
diff --git a/libspu/core/context.cc b/libspu/core/context.cc
new file mode 100644
index 00000000..ff95c166
--- /dev/null
+++ b/libspu/core/context.cc
@@ -0,0 +1,46 @@
+// Copyright 2021 Ant Group Co., Ltd.
+//
+// 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 "libspu/core/context.h"
+
+namespace spu {
+namespace {
+
+std::string genRootObjectId(const std::shared_ptr<yacl::link::Context>& lctx) {
+  // In single-process simulation mode, multi-context need to use different id,
+  // or tracing will not work.
+  if (lctx) {
+    return fmt::format("root-{}", lctx->Rank());
+  }
+
+  return "root";
+}
+
+}  // namespace
+
+SPUContext::SPUContext(const RuntimeConfig& config,
+                       const std::shared_ptr<yacl::link::Context>& lctx)
+    : config_(config),
+      prot_(std::make_unique<Object>(genRootObjectId(lctx))),
+      lctx_(lctx) {}
+
+std::unique_ptr<SPUContext> SPUContext::fork() const {
+  std::shared_ptr<yacl::link::Context> new_lctx =
+      lctx_ ? lctx_->Spawn() : nullptr;
+  auto new_sctx = std::make_unique<SPUContext>(config_, new_lctx);
+  new_sctx->prot_ = prot_->fork();
+  return new_sctx;
+}
+
+}  // namespace spu
diff --git a/libspu/core/context.h b/libspu/core/context.h
new file mode 100644
index 00000000..b887bcf5
--- /dev/null
+++ b/libspu/core/context.h
@@ -0,0 +1,192 @@
+// Copyright 2021 Ant Group Co., Ltd.
+//
+// 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.
+
+#pragma once
+
+#include <memory>
+
+#include "yacl/link/link.h"
+
+#include "libspu/core/object.h"
+#include "libspu/core/prelude.h"
+#include "libspu/core/trace.h"  // TODO: bad reference, but implicitly include too much.
+#include "libspu/core/value.h"
+
+#include "libspu/spu.pb.h"
+
+namespace spu {
+
+// The hal evaluation context for all spu operators.
+class SPUContext final {
+  RuntimeConfig config_;
+
+  // A dynamic object for polymophic(multi-stage) operations.
+  std::unique_ptr<Object> prot_;
+
+  // TODO(jint): do we really need a link here? how about a FHE context.
+  std::shared_ptr<yacl::link::Context> lctx_;
+
+ public:
+  explicit SPUContext(const RuntimeConfig& config,
+                      const std::shared_ptr<yacl::link::Context>& lctx);
+
+  SPUContext(const SPUContext& other) = delete;
+  SPUContext& operator=(const SPUContext& other) = delete;
+  SPUContext(SPUContext&& other) = default;
+
+  // all parties get a 'corresponding' hal context when forked.
+  std::unique_ptr<SPUContext> fork() const;
+
+  //
+  const std::shared_ptr<yacl::link::Context>& lctx() const { return lctx_; }
+
+  // Return current working fixed point fractional bits.
+  size_t getFxpBits() const {
+    const auto fbits = config_.fxp_fraction_bits();
+    SPU_ENFORCE(fbits != 0);
+    return fbits;
+  }
+
+  // Return current working field of MPC engine.
+  FieldType getField() const { return config_.field(); }
+
+  // Return current working runtime config.
+  const RuntimeConfig& config() const { return config_; }
+
+  const std::string& id() { return prot_->id(); }
+  const std::string& pid() { return prot_->pid(); }
+
+  Object* prot() { return prot_.get(); }
+
+  // helper function, forward to caller
+  bool hasKernel(const std::string& name) const {
+    return prot_->hasKernel(name);
+  }
+  Kernel* getKernel(const std::string& name) const {
+    return prot_->getKernel(name);
+  }
+  template <typename StateT>
+  StateT* getState() {
+    return prot_->template getState<StateT>();
+  }
+};
+
+using HalContext [[deprecated("Use SPUContext instead.")]] = SPUContext;
+
+class KernelEvalContext final {
+  // Please keep param types as less as possible.
+  using ParamType = std::variant<  //
+      std::optional<Value>,  // for kernel evaluation failure, give a change
+                             // for the caller to handle failed case.
+      Value,                 // value type
+      Shape,                 //
+      size_t,                // represent size(mmul), shift_bits(shift)
+      bool,                  // binary flag
+      Type,                  // type of type
+      uint128_t              // ring constant
+      >;
+
+  SPUContext* sctx_;
+
+  std::vector<ParamType> params_;
+  ParamType output_;
+
+ public:
+  explicit KernelEvalContext(SPUContext* sctx) : sctx_(sctx) {}
+
+  SPUContext* sctx() { return sctx_; }
+
+  const std::string& id() { return sctx_->id(); }
+  const std::string& pid() { return sctx_->pid(); }
+
+  // helper function, forward to caller.
+  template <typename StateT>
+  StateT* getState() {
+    return sctx_->prot()->template getState<StateT>();
+  }
+
+  // helper function, forward to caller
+  bool hasKernel(const std::string& name) const {
+    return sctx_->prot()->hasKernel(name);
+  }
+
+  size_t numParams() const { return params_.size(); }
+
+  // Steal the output from this evaluation context.
+  //
+  // * usually called by kernel caller.
+  template <typename T = Value>
+  T&& stealOutput() {
+    return std::move(std::get<T>(output_));
+  }
+
+  // Bind an input to this evaluation context.
+  //
+  // * usually called by kernel caller.
+  template <typename T>
+  void bindParam(const T& in) {
+    params_.emplace_back(in);
+  }
+
+  // Get the i'th parameter.
+  //
+  // * usually called by kernel callee.
+  template <typename T>
+  const T& getParam(size_t pos) const {
+    SPU_ENFORCE(pos < params_.size(), "pos={} exceed num of inputs={}", pos,
+                params_.size());
+    return std::get<T>(params_[pos]);
+  }
+
+  // Set the output.
+  //
+  // * usually called by kernel callee.
+  template <typename T = Value>
+  void setOutput(T&& out) {
+    output_ = std::forward<T>(out);
+  }
+};
+
+namespace detail {
+
+template <typename First, typename... Args>
+void bindParams(KernelEvalContext* ectx, First&& head, Args&&... tail) {
+  ectx->bindParam(std::forward<First>(head));
+  if constexpr (sizeof...(Args) > 0) {
+    return bindParams(ectx, std::forward<Args>(tail)...);
+  }
+}
+
+}  // namespace detail
+
+// Dynamic dispath to a kernel according to a symbol name.
+template <typename Ret = Value, typename... Args>
+Ret dynDispatch(SPUContext* sctx, const std::string& name, Args&&... args) {
+  /// Steps of dynamic dispatch.
+  // 1. find a prop kernel.
+  Kernel* kernel = sctx->prot()->getKernel(name);
+
+  // 2. prep parameters (flatten it into an evaluation context).
+  KernelEvalContext ectx(sctx);
+  detail::bindParams(&ectx, std::forward<Args>(args)...);
+
+  // 3. call a visitor, visit a kernel with params.
+  // TODO: use a visitor to call differnet stage of a kernel
+  kernel->evaluate(&ectx);
+
+  // 4. steal the result and return it.
+  return ectx.stealOutput<Ret>();
+}
+
+}  // namespace spu
diff --git a/libspu/core/ndarray_ref.cc b/libspu/core/ndarray_ref.cc
index e4745fdb..01b5af57 100644
--- a/libspu/core/ndarray_ref.cc
+++ b/libspu/core/ndarray_ref.cc
@@ -548,13 +548,11 @@ NdArrayRef unflatten(const ArrayRef& arr, absl::Span<const int64_t> shape) {
                       std::vector<int64_t>(shape.size(), 0), arr.offset());
   }
 
-  // FIXME: due to the current implementation,
-  SPU_ENFORCE(arr.isCompact(), "FIXME: impl assume array is flatten, got {}",
-              arr);
+  ArrayRef compact = arr.isCompact() ? arr : arr.clone();
 
   auto strides = makeCompactStrides(shape);
-  return NdArrayRef(arr.buf(), arr.eltype(), shape, std::move(strides),
-                    arr.offset());
+  return NdArrayRef(compact.buf(), compact.eltype(), shape, std::move(strides),
+                    compact.offset());
 }
 
 namespace {
diff --git a/libspu/core/ndarray_ref.h b/libspu/core/ndarray_ref.h
index e744381a..4ff74905 100644
--- a/libspu/core/ndarray_ref.h
+++ b/libspu/core/ndarray_ref.h
@@ -18,12 +18,12 @@
 #include <vector>
 
 #include "absl/types/span.h"
-#include "spdlog/spdlog.h"
 #include "yacl/base/buffer.h"
 
 #include "libspu/core/array_ref.h"
 #include "libspu/core/shape_util.h"
 #include "libspu/core/type.h"
+#include "libspu/core/vectorize.h"
 
 // #define ITER_DEBUG
 
@@ -87,6 +87,7 @@ class NdArrayRef {
 
   // Return the element type.
   const Type& eltype() const { return eltype_; }
+  Type& eltype() { return eltype_; }
 
   // Return the element size.
   size_t elsize() const { return eltype_.size(); }
@@ -130,11 +131,13 @@ class NdArrayRef {
 
   struct Iterator {
    public:
+    // NOLINTBEGIN, readability-identifier-naming
     using iterator_category = std::forward_iterator_tag;
     using difference_type = std::ptrdiff_t;
     using value_type = std::byte;
     using pointer = std::byte*;
     using reference = std::byte&;
+    // NOLINTEND
 
     explicit Iterator(const NdArrayRef& array, std::vector<int64_t> coord,
                       bool invalid = false)
@@ -292,6 +295,57 @@ class NdArrayRef {
   void eliminate_zero_stride();
 };
 
+template <>
+struct SimdTrait<NdArrayRef> {
+  using Shape = std::vector<int64_t>;  // TODO: use a formal shape definition.
+  using PackInfo = std::vector<Shape>;
+
+  template <typename InputIt>
+  static NdArrayRef pack(InputIt first, InputIt last, PackInfo& pi) {
+    SPU_ENFORCE(first != last);
+
+    int64_t total_numel = 0;
+    const Type ty = first->eltype();
+    for (auto itr = first; itr != last; ++itr) {
+      SPU_ENFORCE(itr->eltype() == ty, "type mismatch {} != {}", itr->eltype(),
+                  ty);
+      total_numel += itr->numel();
+    }
+    NdArrayRef result(first->eltype(), {total_numel});
+    int64_t offset = 0;
+    for (; first != last; ++first) {
+      NdArrayRef slice(result.buf(), ty, first->shape(),
+                       makeCompactStrides(first->shape()), offset);
+      const std::vector<int64_t> start_index(first->ndim(), 0);
+      slice.copy_slice(*first, start_index, start_index, first->numel());
+      pi.push_back(first->shape());
+      offset += first->numel() * ty.size();
+    }
+    return result;
+  }
+
+  template <typename OutputIt>
+  static OutputIt unpack(const NdArrayRef& v, OutputIt result,
+                         const PackInfo& pi) {
+    int64_t total_num = 0;
+    for (const auto& shape : pi) {
+      total_num += calcNumel(shape);
+    }
+
+    SPU_ENFORCE(v.numel() == total_num, "split number mismatch {} != {}",
+                v.numel(), total_num);
+
+    int64_t offset = 0;
+    for (const auto& shape : pi) {
+      *result++ = NdArrayRef(v.buf(), v.eltype(), shape,
+                             makeCompactStrides(shape), offset);
+      offset += calcNumel(shape) * v.elsize();
+    }
+
+    return result;
+  }
+};
+
 // Unflatten a 1d-array to an ndarray.
 NdArrayRef unflatten(const ArrayRef& arr, absl::Span<const int64_t> shape);
 
diff --git a/libspu/mpc/object.cc b/libspu/core/object.cc
similarity index 81%
rename from libspu/mpc/object.cc
rename to libspu/core/object.cc
index a56e4a59..127d4756 100644
--- a/libspu/mpc/object.cc
+++ b/libspu/core/object.cc
@@ -12,15 +12,15 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "libspu/mpc/object.h"
+#include "libspu/core/object.h"
 
-namespace spu::mpc {
+namespace spu {
 
 std::unique_ptr<State> State::fork() {
   SPU_THROW("Not implemented, the sub class should override this");
 }
 
-std::unique_ptr<Object> Object::fork() {
+std::unique_ptr<Object> Object::fork() const {
   auto new_id = fmt::format("{}-{}", id_, child_counter_++);
   auto new_obj = std::make_unique<Object>(new_id, id_);
   new_obj->kernels_ = kernels_;
@@ -39,21 +39,22 @@ bool Object::hasLowCostFork() const {
   return true;
 }
 
-void Object::regKernel(std::string_view name, std::unique_ptr<Kernel> kernel) {
+void Object::regKernel(const std::string& name,
+                       std::unique_ptr<Kernel> kernel) {
   const auto itr = kernels_.find(name);
   SPU_ENFORCE(itr == kernels_.end(), "kernel={} already exist", name);
   kernels_.insert({name, std::move(kernel)});
 }
 
-Kernel* Object::getKernel(std::string_view name) {
+Kernel* Object::getKernel(const std::string& name) const {
   const auto itr = kernels_.find(name);
   SPU_ENFORCE(itr != kernels_.end(), "kernel={} not found", name);
   return itr->second.get();
 }
 
-bool Object::hasKernel(std::string_view name) const {
+bool Object::hasKernel(const std::string& name) const {
   const auto itr = kernels_.find(name);
   return itr != kernels_.end();
 }
 
-}  // namespace spu::mpc
+}  // namespace spu
diff --git a/libspu/core/object.h b/libspu/core/object.h
new file mode 100644
index 00000000..d462b951
--- /dev/null
+++ b/libspu/core/object.h
@@ -0,0 +1,159 @@
+// Copyright 2021 Ant Group Co., Ltd.
+//
+// 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.
+
+#pragma once
+
+#include <map>
+#include <memory>
+#include <optional>
+#include <string>
+#include <variant>
+
+#include "libspu/core/array_ref.h"
+#include "libspu/core/cexpr.h"
+#include "libspu/core/prelude.h"
+
+namespace spu {
+
+// TODO(jint) document the object model.
+// - KernelEvalContext: the calling convention
+// - Kernel: the dynamic member function.
+// - State: the dynamic member variable.
+// - Object: the dynamic binding object.
+
+class KernelEvalContext;
+class Kernel {
+ public:
+  enum class Kind {
+    // Indicate the kernel's complexity is static known.
+    //
+    // Typically, static kernel does not depend on runtime options, such like
+    // selecting different kernels according to different configs.
+    //
+    // By default, we should make kernels as 'Static' as possible.
+    Static,
+
+    // Indicate the kernel depends on runtime options, this kind of kernel is
+    // hard to analysis statically.
+    Dynamic,
+  };
+
+  virtual ~Kernel() = default;
+
+  virtual Kind kind() const { return Kind::Static; }
+
+  // Calculate number of comm rounds required for this kernel.
+  virtual ce::CExpr latency() const { return nullptr; }
+
+  // Calculate number of comm in bits.
+  virtual ce::CExpr comm() const { return nullptr; }
+
+  // Some kernel's communication can not be measured/implemented easily, this
+  // field tells the tolerance of theory value and implementation diff, in
+  // percentage. i.e.
+  //
+  // kernel's cost is infinitely close to 2.
+  //   comm(kernel) = 1 + 1/2 + 1/4 + 1/8 + ... = lim(2)
+  //
+  // in implementation, program may not handle bit, but instead use byte as the
+  // minimum unit, this make the error even larger.
+  virtual float getCommTolerance() const { return 0.0; }
+
+  // Evaluate this protocol within given context.
+  virtual void evaluate(KernelEvalContext* ectx) const = 0;
+};
+
+class State {
+ public:
+  virtual ~State() = default;
+
+  // Return true if the state could be forked with low cost.
+  virtual bool hasLowCostFork() const { return false; }
+
+  // TODO: this is a const method.
+  virtual std::unique_ptr<State> fork();
+};
+
+// A dynamic object contains a set of kernels and a set of states.
+//
+// Class that inherit from this class could do `dynamic binding`.
+class Object final {
+  std::map<std::string, std::shared_ptr<Kernel>> kernels_;
+  std::map<std::string, std::unique_ptr<State>> states_;
+
+  std::string id_;   // this object id.
+  std::string pid_;  // parent id.
+
+  mutable int64_t child_counter_ = 0;
+
+ public:
+  explicit Object(std::string id, std::string pid = "")
+      : id_(std::move(id)), pid_(std::move(pid)) {}
+
+  virtual ~Object() = default;
+
+  const std::string& id() const { return id_; }
+  const std::string& pid() const { return pid_; }
+
+  //
+  std::unique_ptr<Object> fork() const;
+
+  bool hasLowCostFork() const;
+
+  void regKernel(const std::string& name, std::unique_ptr<Kernel> kernel);
+
+  template <typename KernelT>
+  void regKernel() {
+    regKernel(KernelT::kBindName, std::make_unique<KernelT>());
+  }
+
+  template <typename KernelT>
+  void regKernel(const std::string& name) {
+    return regKernel(name, std::make_unique<KernelT>());
+  }
+
+  Kernel* getKernel(const std::string& name) const;
+  bool hasKernel(const std::string& name) const;
+
+  void addState(const std::string& name, std::unique_ptr<State> state) {
+    const auto& itr = states_.find(name);
+    SPU_ENFORCE(itr == states_.end(), "state={} already exist", name);
+    states_.emplace(name, std::move(state));
+  }
+
+  template <typename StateT, typename... Args>
+  void addState(Args&&... args) {
+    addState(StateT::kBindName,
+             std::make_unique<StateT>(std::forward<Args>(args)...));
+  }
+
+  template <typename StateT>
+  StateT* getState() {
+    const auto& itr = states_.find(StateT::kBindName);
+    SPU_ENFORCE(itr != states_.end(), "state={} not found", StateT::kBindName);
+    return dynamic_cast<StateT*>(itr->second.get());
+  }
+
+  //
+  std::vector<std::string> getKernelNames() const {
+    std::vector<std::string> names;
+    names.reserve(kernels_.size());
+    for (auto const& itr : kernels_) {
+      names.push_back(itr.first);
+    }
+    return names;
+  }
+};
+
+}  // namespace spu
diff --git a/libspu/core/shape_util.h b/libspu/core/shape_util.h
index b982e8db..86c6a5fa 100644
--- a/libspu/core/shape_util.h
+++ b/libspu/core/shape_util.h
@@ -71,4 +71,33 @@ inline size_t calcFlattenOffset(absl::Span<const int64_t> indices,
   return offset;
 }
 
+using ShapeView = absl::Span<int64_t const>;
+
+class Shape : public std::vector<int64_t> {
+  using Base = std::vector<int64_t>;
+
+ public:
+  Shape() = default;
+  Shape(const Shape &other) = default;
+  Shape &operator=(const Shape &other) = default;
+
+  Shape(std::initializer_list<int64_t> list) : Base(list) {}
+  Shape(iterator begin, iterator end) : Base(begin, end) {}
+  explicit Shape(size_t size, int64_t element = 0) : Base(size, element) {}
+  explicit Shape(const std::vector<int64_t> &other) : Base(other) {}
+  explicit Shape(absl::Span<int64_t const> other)
+      : Base(other.begin(), other.end()) {}
+
+  int64_t ndim() const { return static_cast<int64_t>(size()); }
+
+  int64_t dim(int64_t idx) const {
+    SPU_ENFORCE(idx < ndim(), "invalid idx={}, ndim={}", idx, ndim());
+    return at(idx);
+  }
+
+  int64_t numel() const {
+    return calcNumel(absl::MakeConstSpan(data(), size()));
+  }
+};
+
 }  // namespace spu
diff --git a/libspu/core/trace.h b/libspu/core/trace.h
index 73fcca99..0f0efebf 100644
--- a/libspu/core/trace.h
+++ b/libspu/core/trace.h
@@ -351,13 +351,13 @@ std::shared_ptr<Tracer> getTracer(const std::string& id,
 
 // trace an mpc layer dispatch
 #define SPU_TRACE_MPC_DISP(CTX, ...)                                   \
-  SPU_TRACE_ACTION(GET_TRACER(CTX->caller()), (TR_MPC | TR_LOG), (~0), \
-                   kBindName, ##__VA_ARGS__)
+  SPU_TRACE_ACTION(GET_TRACER(CTX), (TR_MPC | TR_LOG), (~0), __func__, \
+                   ##__VA_ARGS__)
 
 // trace an mpc layer leaf
 #define SPU_TRACE_MPC_LEAF(CTX, ...)                                        \
-  SPU_TRACE_ACTION(GET_TRACER(CTX->caller()), (TR_MPC | TR_LAR), (~TR_MPC), \
-                   kBindName, ##__VA_ARGS__)
+  SPU_TRACE_ACTION(GET_TRACER(CTX), (TR_MPC | TR_LAR), (~TR_MPC), __func__, \
+                   ##__VA_ARGS__)
 
 // Debug purpose only.
 class MemProfilingGuard {
diff --git a/libspu/kernel/value.cc b/libspu/core/value.cc
similarity index 90%
rename from libspu/kernel/value.cc
rename to libspu/core/value.cc
index 74c08b7d..56e02882 100644
--- a/libspu/kernel/value.cc
+++ b/libspu/core/value.cc
@@ -12,7 +12,7 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "libspu/kernel/value.h"
+#include "libspu/core/value.h"
 
 #include <algorithm>
 #include <cstdint>
@@ -122,4 +122,14 @@ std::ostream& operator<<(std::ostream& out, const Value& v) {
   return out;
 }
 
+std::tuple<ArrayRef, Shape, DataType> UnwrapValue(const Value& val) {
+  return std::make_tuple(flatten(val.data()), Shape(val.shape()), val.dtype());
+}
+
+Value WrapValue(const ArrayRef& arr, absl::Span<int64_t const> shape,
+                DataType dtype) {
+  auto ndarr = unflatten(arr, shape);
+  return Value(ndarr, dtype);
+}
+
 }  // namespace spu
diff --git a/libspu/kernel/value.h b/libspu/core/value.h
similarity index 51%
rename from libspu/kernel/value.h
rename to libspu/core/value.h
index 03647445..410d322f 100644
--- a/libspu/kernel/value.h
+++ b/libspu/core/value.h
@@ -21,6 +21,7 @@
 #include "libspu/core/ndarray_ref.h"
 #include "libspu/core/type.h"
 #include "libspu/core/type_util.h"
+#include "libspu/core/vectorize.h"
 
 namespace spu {
 
@@ -40,6 +41,7 @@ class Value final {
 
   // Get the concrete storage type.
   const Type& storage_type() const { return data_.eltype(); }
+  Type& storage_type() { return data_.eltype(); }
 
   // Access the underline storage data.
   const NdArrayRef& data() const { return data_; }
@@ -72,6 +74,65 @@ class Value final {
   Value clone() const;
 };
 
+// Helper function to legacy kernels. TODO: drop this wraps.
+std::tuple<ArrayRef, Shape, DataType> UnwrapValue(const Value& val);
+Value WrapValue(const ArrayRef& arr, absl::Span<int64_t const> shape,
+                DataType dtype = DT_INVALID);
+
+template <>
+struct SimdTrait<Value> {
+  using Shape = std::vector<int64_t>;  // TODO: use a formal shape definition.
+  using PackInfo = std::vector<Shape>;
+
+  template <typename InputIt>
+  static Value pack(InputIt first, InputIt last, PackInfo& pi) {
+    SPU_ENFORCE(first != last);
+
+    int64_t total_numel = 0;
+    const Type ty = first->storage_type();
+    const auto dtype = first->dtype();
+    for (auto itr = first; itr != last; ++itr) {
+      SPU_ENFORCE(itr->storage_type() == ty, "type mismatch {} != {}",
+                  itr->storage_type(), ty);
+      SPU_ENFORCE(itr->dtype() == dtype, "dtype mismatch {} != {}",
+                  itr->dtype(), dtype);
+      total_numel += itr->numel();
+    }
+    NdArrayRef result(ty, {total_numel});
+    int64_t offset = 0;
+    for (; first != last; ++first) {
+      NdArrayRef slice(result.buf(), ty, first->shape(),
+                       makeCompactStrides(first->shape()), offset);
+      const std::vector<int64_t> start_index(first->shape().size(), 0);
+      slice.copy_slice(first->data(), start_index, start_index, first->numel());
+      pi.push_back(first->shape());
+      offset += first->numel() * ty.size();
+    }
+    return Value(result, dtype);
+  }
+
+  template <typename OutputIt>
+  static OutputIt unpack(const Value& v, OutputIt result, const PackInfo& pi) {
+    int64_t total_num = 0;
+    for (const auto& shape : pi) {
+      total_num += calcNumel(shape);
+    }
+
+    SPU_ENFORCE(v.numel() == total_num, "split number mismatch {} != {}",
+                v.numel(), total_num);
+
+    int64_t offset = 0;
+    for (const auto& shape : pi) {
+      auto arr = NdArrayRef(v.data().buf(), v.storage_type(), shape,
+                            makeCompactStrides(shape), offset);
+      *result++ = Value(arr, v.dtype());
+      offset += calcNumel(shape) * v.elsize();
+    }
+
+    return result;
+  }
+};
+
 std::ostream& operator<<(std::ostream& out, const Value& v);
 
 }  // namespace spu
diff --git a/libspu/device/BUILD.bazel b/libspu/device/BUILD.bazel
index bc4800a5..671b6ce7 100644
--- a/libspu/device/BUILD.bazel
+++ b/libspu/device/BUILD.bazel
@@ -47,9 +47,9 @@ spu_cc_library(
         ":device_cc_proto",
         ":symbol_table",
         "//libspu:spu_cc_proto",
+        "//libspu/core:context",
         "//libspu/core:pt_buffer_view",
-        "//libspu/kernel:context",
-        "//libspu/kernel:value",
+        "//libspu/core:value",
         "//libspu/kernel/hal:constants",
         "//libspu/kernel/hal:public_helper",
         "//libspu/mpc:factory",
@@ -73,9 +73,9 @@ spu_cc_library(
         ":device_cc_proto",
         ":symbol_table",
         "//libspu:spu_cc_proto",
+        "//libspu/core:context",
+        "//libspu/core:value",
         "//libspu/dialect:pphlo_dialect",
-        "//libspu/kernel:context",
-        "//libspu/kernel:value",
         "@llvm-project//mlir:IR",
     ],
 )
@@ -113,6 +113,6 @@ spu_cc_library(
     deps = [
         ":device_cc_proto",
         "//libspu/core:pt_buffer_view",
-        "//libspu/kernel:value",
+        "//libspu/core:value",
     ],
 )
diff --git a/libspu/device/api.cc b/libspu/device/api.cc
index b5ecc644..d8892820 100644
--- a/libspu/device/api.cc
+++ b/libspu/device/api.cc
@@ -124,7 +124,7 @@ void takeSnapshot(size_t rank, const RuntimeConfig &rt_config,
   dump_file << snapshot.SerializeAsString();
 }
 
-void printProfilingData(spu::HalContext *hctx, const std::string &name,
+void printProfilingData(spu::SPUContext *sctx, const std::string &name,
                         const ExecutionStats &exec_stats,
                         const CommunicationStats &comm_stats) {
   // print overall information
@@ -139,7 +139,7 @@ void printProfilingData(spu::HalContext *hctx, const std::string &name,
   {
     std::map<ActionKey, ActionStats> stats;
 
-    const auto &tracer = GET_TRACER(hctx);
+    const auto &tracer = GET_TRACER(sctx);
     const auto &records = tracer->getProfState()->getRecords();
 
     for (const auto &rec : records) {
@@ -176,7 +176,7 @@ void printProfilingData(spu::HalContext *hctx, const std::string &name,
               comm_stats.send_bytes, comm_stats.send_actions);
 }
 
-void setupTrace(spu::HalContext *hctx, const spu::RuntimeConfig &rt_config) {
+void setupTrace(spu::SPUContext *sctx, const spu::RuntimeConfig &rt_config) {
   int64_t tr_flag = 0;
   // TODO: Support tracing for parallel op execution
   if (rt_config.enable_action_trace() &&
@@ -194,8 +194,8 @@ void setupTrace(spu::HalContext *hctx, const spu::RuntimeConfig &rt_config) {
     tr_flag |= TR_REC;
   }
 
-  initTrace(hctx->id(), tr_flag);
-  GET_TRACER(hctx)->getProfState()->clearRecords();
+  initTrace(sctx->id(), tr_flag);
+  GET_TRACER(sctx)->getProfState()->clearRecords();
 }
 
 void SPUErrorHandler(void *use_data, const char *reason, bool gen_crash_diag) {
@@ -218,13 +218,13 @@ void installLLVMErrorHandler() {
 
 }  // namespace
 
-void executeImpl(OpExecutor *executor, spu::HalContext *hctx,
+void executeImpl(OpExecutor *executor, spu::SPUContext *sctx,
                  const ExecutableProto &executable, SymbolTable *env) {
-  setupTrace(hctx, hctx->rt_config());
+  setupTrace(sctx, sctx->config());
   installLLVMErrorHandler();
 
   CommunicationStats comm_stats;
-  comm_stats.reset(hctx->lctx());
+  comm_stats.reset(sctx->lctx());
   ExecutionStats exec_stats;
 
   // prepare inputs from environment.
@@ -238,10 +238,10 @@ void executeImpl(OpExecutor *executor, spu::HalContext *hctx,
   }
 
   // TODO: rename this flag, enable_execution_dump?
-  const RuntimeConfig rt_config = hctx->rt_config();
+  const RuntimeConfig rt_config = sctx->config();
   if (rt_config.enable_processor_dump()) {
-    const bool isRefHal = hctx->lctx() == nullptr;
-    const size_t rank = isRefHal ? 0 : hctx->lctx()->Rank();
+    const bool isRefHal = sctx->lctx() == nullptr;
+    const size_t rank = isRefHal ? 0 : sctx->lctx()->Rank();
     takeSnapshot(rank, rt_config, executable, *env);
   }
 
@@ -274,7 +274,7 @@ void executeImpl(OpExecutor *executor, spu::HalContext *hctx,
       mlir_ctx.enableMultithreading();
       mlir_ctx.enterMultiThreadedExecution();
     }
-    outputs = runRegion(executor, hctx, nullptr, entry_function.getBody(),
+    outputs = runRegion(executor, sctx, nullptr, entry_function.getBody(),
                         inputs, opts);
 
     if (opts.do_parallel) {
@@ -290,18 +290,18 @@ void executeImpl(OpExecutor *executor, spu::HalContext *hctx,
     }
   }
 
-  comm_stats.diff(hctx->lctx());
-  if ((getGlobalTraceFlag(hctx->id()) & TR_REC) != 0) {
-    printProfilingData(hctx, executable.name(), exec_stats, comm_stats);
+  comm_stats.diff(sctx->lctx());
+  if ((getGlobalTraceFlag(sctx->id()) & TR_REC) != 0) {
+    printProfilingData(sctx, executable.name(), exec_stats, comm_stats);
   }
 }
 
-void execute(OpExecutor *executor, spu::HalContext *hctx,
+void execute(OpExecutor *executor, spu::SPUContext *sctx,
              const spu::ExecutableProto &executable, SymbolTable *env) {
-  return executeImpl(executor, hctx, executable, env);
+  return executeImpl(executor, sctx, executable, env);
 }
 
-void execute(OpExecutor *executor, spu::HalContext *hctx,
+void execute(OpExecutor *executor, spu::SPUContext *sctx,
              const std::string &text,
              const std::vector<std::string> &input_names,
              const std::vector<std::string> &output_names, SymbolTable *env) {
@@ -312,7 +312,7 @@ void execute(OpExecutor *executor, spu::HalContext *hctx,
                                         output_names.end()};
   executable.set_code(text);
 
-  return executeImpl(executor, hctx, executable, env);
+  return executeImpl(executor, sctx, executable, env);
 }
 
 }  // namespace spu::device
diff --git a/libspu/device/api.h b/libspu/device/api.h
index 1da01186..215390aa 100644
--- a/libspu/device/api.h
+++ b/libspu/device/api.h
@@ -17,20 +17,20 @@
 #include <string>
 #include <vector>
 
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 #include "libspu/device/executor.h"
 #include "libspu/device/symbol_table.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
 
 #include "libspu/spu.pb.h"
 
 namespace spu::device {
 
-void execute(OpExecutor *executor, HalContext *hctx,
+void execute(OpExecutor *executor, SPUContext *sctx,
              const ExecutableProto &executable, SymbolTable *env);
 
 ///
-void execute(OpExecutor *executor, spu::HalContext *hctx,
+void execute(OpExecutor *executor, spu::SPUContext *sctx,
              const std::string &text,
              const std::vector<std::string> &input_names,
              const std::vector<std::string> &output_names, SymbolTable *env);
diff --git a/libspu/device/executor.cc b/libspu/device/executor.cc
index 41833493..f8ed0f10 100644
--- a/libspu/device/executor.cc
+++ b/libspu/device/executor.cc
@@ -24,9 +24,9 @@
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/Value.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/prelude.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/value.h"
 
 namespace spu::device {
 
@@ -98,7 +98,7 @@ void SymbolScope::addValue(mlir::Value key, spu::Value &&val) {
 }
 
 std::vector<spu::Value> runRegion(OpExecutor *executor,                 //
-                                  HalContext *hctx,                     //
+                                  SPUContext *sctx,                     //
                                   SymbolScope *parent_scope,            //
                                   mlir::Region &region,                 //
                                   absl::Span<spu::Value const> params,  //
@@ -117,19 +117,19 @@ std::vector<spu::Value> runRegion(OpExecutor *executor,                 //
 
   SPU_ENFORCE(region.hasOneBlock());
   if (opts.do_parallel) {
-    return runBlockParallel(executor, hctx, &sscope, region.front(), params,
+    return runBlockParallel(executor, sctx, &sscope, region.front(), params,
                             opts);
   } else {
-    return runBlock(executor, hctx, &sscope, region.front(), params, opts);
+    return runBlock(executor, sctx, &sscope, region.front(), params, opts);
   }
 }
 
-std::vector<spu::Value> runBlock(OpExecutor *executor, HalContext *hctx,
+std::vector<spu::Value> runBlock(OpExecutor *executor, SPUContext *sctx,
                                  SymbolScope *symbols, mlir::Block &block,
                                  absl::Span<spu::Value const> params,
                                  const ExecutionOptions &opts) {
   for (auto &op : block.without_terminator()) {
-    executor->runKernel(hctx, symbols, op, opts);
+    executor->runKernel(sctx, symbols, op, opts);
   }
 
   if (auto *termOp = block.getTerminator()) {
@@ -152,7 +152,7 @@ struct SymbolTableEvent {
 };
 
 class OpExecTask final {
-  std::unique_ptr<HalContext> hctx_ = nullptr;
+  std::unique_ptr<SPUContext> sctx_ = nullptr;
   // here we assume executor is thread-safe (stateless)
   OpExecutor *executor_ = nullptr;
   SymbolScope *sscope_ = nullptr;
@@ -162,10 +162,10 @@ class OpExecTask final {
 
  public:
   OpExecTask() = default;
-  explicit OpExecTask(std::unique_ptr<HalContext> hctx, OpExecutor *executor,
+  explicit OpExecTask(std::unique_ptr<SPUContext> sctx, OpExecutor *executor,
                       SymbolScope *sscope, mlir::Operation *op,
                       SymbolTableEvent *event)
-      : hctx_(std::move(hctx)),
+      : sctx_(std::move(sctx)),
         executor_(executor),
         sscope_(sscope),
         op_(op),
@@ -199,13 +199,13 @@ class OpExecTask final {
       event_->cv.wait(lk, [this] { return ready(); });
     }
 
-    executor_->runKernel(hctx_.get(), sscope_, *op_);
+    executor_->runKernel(sctx_.get(), sscope_, *op_);
     std::unique_lock lk(event_->mutex);
     event_->cv.notify_all();
   }
 };
 
-std::vector<spu::Value> runBlockParallel(OpExecutor *executor, HalContext *hctx,
+std::vector<spu::Value> runBlockParallel(OpExecutor *executor, SPUContext *sctx,
                                          SymbolScope *symbols,
                                          mlir::Block &block,
                                          absl::Span<spu::Value const> params,
@@ -217,7 +217,7 @@ std::vector<spu::Value> runBlockParallel(OpExecutor *executor, HalContext *hctx,
   std::vector<OpExecTask> tasks;
   std::vector<std::future<void>> futures;
   for (auto &op : block.without_terminator()) {
-    tasks.emplace_back(hctx->fork(), executor, symbols, &op, &st_event);
+    tasks.emplace_back(sctx->fork(), executor, symbols, &op, &st_event);
   }
 
   futures.reserve(tasks.size());
diff --git a/libspu/device/executor.h b/libspu/device/executor.h
index 6cab395d..9c0e21d1 100644
--- a/libspu/device/executor.h
+++ b/libspu/device/executor.h
@@ -21,8 +21,8 @@
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/MLIRContext.h"
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::device {
 
@@ -71,28 +71,28 @@ class OpExecutor {
   virtual bool hasKernel(mlir::Operation &op) const = 0;
 
   // run a kernel in a given region.
-  virtual void runKernelImpl(HalContext *hctx, SymbolScope *sscope,
+  virtual void runKernelImpl(SPUContext *sctx, SymbolScope *sscope,
                              mlir::Operation &op,
                              const ExecutionOptions &opts) = 0;
 
-  void runKernel(HalContext *hctx, SymbolScope *sscope, mlir::Operation &op,
+  void runKernel(SPUContext *sctx, SymbolScope *sscope, mlir::Operation &op,
                  const ExecutionOptions &opts = {}) {
-    return runKernelImpl(hctx, sscope, op, opts);
+    return runKernelImpl(sctx, sscope, op, opts);
   }
 };
 
-std::vector<spu::Value> runRegion(OpExecutor *executor, HalContext *hctx,
+std::vector<spu::Value> runRegion(OpExecutor *executor, SPUContext *sctx,
                                   SymbolScope *parent_scope,
                                   mlir::Region &region,
                                   absl::Span<spu::Value const> params,
                                   const ExecutionOptions &opts = {});
 
-std::vector<spu::Value> runBlock(OpExecutor *executor, HalContext *hctx,
+std::vector<spu::Value> runBlock(OpExecutor *executor, SPUContext *sctx,
                                  SymbolScope *symbols, mlir::Block &block,
                                  absl::Span<spu::Value const> params,
                                  const ExecutionOptions &opts);
 
-std::vector<spu::Value> runBlockParallel(OpExecutor *executor, HalContext *hctx,
+std::vector<spu::Value> runBlockParallel(OpExecutor *executor, SPUContext *sctx,
                                          SymbolScope *symbols,
                                          mlir::Block &block,
                                          absl::Span<spu::Value const> params,
diff --git a/libspu/device/io.cc b/libspu/device/io.cc
index a1f2a4b5..d4e356be 100644
--- a/libspu/device/io.cc
+++ b/libspu/device/io.cc
@@ -102,7 +102,7 @@ NdArrayRef IoClient::combineShares(absl::Span<spu::Value const> values) {
   return decodeFromRing(encoded, dtype, fxp_bits);
 }
 
-ColocatedIo::ColocatedIo(HalContext *hctx) : hctx_(hctx) {}
+ColocatedIo::ColocatedIo(SPUContext *sctx) : sctx_(sctx) {}
 
 void ColocatedIo::hostSetVar(const std::string &name, const PtBufferView &bv,
                              Visibility vtype) {
@@ -118,7 +118,7 @@ NdArrayRef ColocatedIo::hostGetVar(const std::string &name) const {
   const spu::Value &v = symbols_.getVar(name);
 
   if (v.isPublic()) {
-    return kernel::hal::dump_public(hctx_, v);
+    return kernel::hal::dump_public(sctx_, v);
   } else if (v.isSecret()) {
     SPU_THROW("not implemented");
     // TODO: test the secret's owner is self,
@@ -199,9 +199,9 @@ void ColocatedIo::sync() {
   // Currently, we implement the naive method, that is, using hal link context
   // to send all shares to others.
 
-  const auto &lctx = hctx_->lctx();
+  const auto &lctx = sctx_->lctx();
 
-  IoClient io(lctx->WorldSize(), hctx_->rt_config());
+  IoClient io(lctx->WorldSize(), sctx_->config());
   std::vector<SymbolTableProto> shares_per_party(lctx->WorldSize());
   for (const auto &[name, priv] : unsynced_) {
     const auto &arr = priv.arr;
diff --git a/libspu/device/io.h b/libspu/device/io.h
index d0a7b841..7e810feb 100644
--- a/libspu/device/io.h
+++ b/libspu/device/io.h
@@ -18,10 +18,10 @@
 #include <memory>
 #include <vector>
 
+#include "libspu/core/context.h"
 #include "libspu/core/pt_buffer_view.h"
+#include "libspu/core/value.h"
 #include "libspu/device/symbol_table.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
 #include "libspu/mpc/io_interface.h"
 
 namespace spu::device {
@@ -115,7 +115,7 @@ class IoClient {
 
 class ColocatedIo {
   // The hardware context.
-  HalContext *hctx_;
+  SPUContext *sctx_;
 
   // the place that variables will be store/load.
   SymbolTable symbols_;
@@ -128,11 +128,11 @@ class ColocatedIo {
   std::map<std::string, PrivData> unsynced_;
 
  public:
-  explicit ColocatedIo(HalContext *hctx);
+  explicit ColocatedIo(SPUContext *sctx);
 
-  size_t getWorldSize() const { return hctx_->lctx()->WorldSize(); }
+  size_t getWorldSize() const { return sctx_->lctx()->WorldSize(); }
 
-  size_t getRank() const { return hctx_->lctx()->Rank(); }
+  size_t getRank() const { return sctx_->lctx()->Rank(); }
 
   void hostSetVar(const std::string &name, const PtBufferView &bv,
                   Visibility vtype = VIS_SECRET);
diff --git a/libspu/device/pphlo/BUILD.bazel b/libspu/device/pphlo/BUILD.bazel
index 34757e67..ae6b6557 100644
--- a/libspu/device/pphlo/BUILD.bazel
+++ b/libspu/device/pphlo/BUILD.bazel
@@ -37,6 +37,7 @@ spu_cc_test(
         "//libspu/device:api",
         "//libspu/device:io",
         "//libspu/device:test_utils",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -45,9 +46,9 @@ spu_cc_library(
     srcs = ["pphlo_verifier.cc"],
     hdrs = ["pphlo_verifier.h"],
     deps = [
+        "//libspu/core:context",
+        "//libspu/core:value",
         "//libspu/dialect:pphlo_dialect",
-        "//libspu/kernel:context",
-        "//libspu/kernel:value",
         "//libspu/kernel/hlo",
         "@stablehlo//:reference_ops",
     ],
@@ -59,6 +60,7 @@ spu_cc_test(
     deps = [
         ":pphlo_verifier",
         "//libspu/device:test_utils",
+        "//libspu/kernel:test_util",
         "//libspu/mpc/utils:simulate",
         "@llvm-project//mlir:Parser",
     ],
diff --git a/libspu/device/pphlo/executor_debug_runner.cc b/libspu/device/pphlo/executor_debug_runner.cc
index a81a6ebb..3876deed 100644
--- a/libspu/device/pphlo/executor_debug_runner.cc
+++ b/libspu/device/pphlo/executor_debug_runner.cc
@@ -21,12 +21,12 @@
 #include "llvm/Support/CommandLine.h"
 #include "spdlog/spdlog.h"
 
+#include "libspu/core/value.h"
 #include "libspu/device/api.h"
 #include "libspu/device/pphlo/pphlo_executor.h"
 #include "libspu/device/symbol_table.h"
 #include "libspu/device/test_utils.h"
 #include "libspu/kernel/hal/debug.h"
-#include "libspu/kernel/value.h"
 #include "libspu/mpc/utils/simulate.h"
 
 llvm::cl::opt<std::string> SnapshotDir(
@@ -64,11 +64,11 @@ std::shared_ptr<yacl::link::Context> MakeLink(const std::string &parties,
   return lctx;
 }
 
-std::unique_ptr<spu::HalContext> MakeHalContext(
+std::unique_ptr<spu::SPUContext> MakeSPUContext(
     const spu::device::SnapshotProto &snapshot) {
   auto lctx = MakeLink(Parties.getValue(), Rank.getValue());
 
-  return std::make_unique<spu::HalContext>(snapshot.runtime_cfg(), lctx);
+  return std::make_unique<spu::SPUContext>(snapshot.runtime_cfg(), lctx);
 }
 
 spu::device::SnapshotProto ParseSnapshotFile(
@@ -92,16 +92,16 @@ void RpcBasedRunner(const std::filesystem::path &snapshot_dir) {
   auto snapshot_file =
       snapshot_dir / fmt::format("snapshot_{}.spu", Rank.getValue());
   spu::device::SnapshotProto snapshot = ParseSnapshotFile(snapshot_file);
-  auto hctx = MakeHalContext(snapshot);
+  auto sctx = MakeSPUContext(snapshot);
 
   spu::device::SymbolTable table =
       spu::device::SymbolTable::fromProto(snapshot.environ());
 
   spu::device::pphlo::PPHloExecutor executor;
 
-  SPDLOG_INFO("Run with config {}", hctx->rt_config().DebugString());
+  SPDLOG_INFO("Run with config {}", sctx->config().DebugString());
 
-  spu::device::execute(&executor, hctx.get(), snapshot.executable(), &table);
+  spu::device::execute(&executor, sctx.get(), snapshot.executable(), &table);
 }
 
 void MemBasedRunner(const std::filesystem::path &snapshot_dir) {
@@ -114,12 +114,12 @@ void MemBasedRunner(const std::filesystem::path &snapshot_dir) {
 
         spu::device::SnapshotProto snapshot = ParseSnapshotFile(snapshot_file);
 
-        spu::HalContext hctx(snapshot.runtime_cfg(), lctx);
+        spu::SPUContext sctx(snapshot.runtime_cfg(), lctx);
 
         spu::device::pphlo::PPHloExecutor executor;
         spu::device::SymbolTable table =
             spu::device::SymbolTable::fromProto(snapshot.environ());
-        spu::device::execute(&executor, &hctx, snapshot.executable(), &table);
+        spu::device::execute(&executor, &sctx, snapshot.executable(), &table);
       });
 }
 
diff --git a/libspu/device/pphlo/pphlo_executor.cc b/libspu/device/pphlo/pphlo_executor.cc
index 588c6ac9..a0558f1e 100644
--- a/libspu/device/pphlo/pphlo_executor.cc
+++ b/libspu/device/pphlo/pphlo_executor.cc
@@ -199,10 +199,10 @@ void addValue(SymbolScope *scope, mlir::Value key, spu::Value &&val,
 
 //
 #define STANDARD_UNARY_OP_EXEC_IMPL(OpName, KernelName)                     \
-  void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope, \
+  void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope, \
                mlir::pphlo::OpName &op, const ExecutionOptions &opts) {     \
     const auto in = lookupValue(sscope, op.getOperand(), opts);             \
-    auto ret = kernel::hlo::KernelName(hctx, in);                           \
+    auto ret = kernel::hlo::KernelName(sctx, in);                           \
     addValue(sscope, op.getResult(), std::move(ret), opts);                 \
   }
 
@@ -225,11 +225,11 @@ STANDARD_UNARY_OP_EXEC_IMPL(RoundOp, Round_AFZ)
 #undef STANDARD_UNARY_OP_EXEC_IMPL
 
 #define STANDARD_BINARY_OP_EXEC_IMPL(OpName, KernelName)                      \
-  void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,   \
+  void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,   \
                mlir::pphlo::OpName &op, const ExecutionOptions &opts) {       \
     addValue(                                                                 \
         sscope, op.getResult(),                                               \
-        kernel::hlo::KernelName(hctx, lookupValue(sscope, op.getLhs(), opts), \
+        kernel::hlo::KernelName(sctx, lookupValue(sscope, op.getLhs(), opts), \
                                 lookupValue(sscope, op.getRhs(), opts)),      \
         opts);                                                                \
   }
@@ -255,7 +255,7 @@ STANDARD_BINARY_OP_EXEC_IMPL(ShiftRightLogicalOp, Rshift)
 
 #undef STANDARD_BINARY_OP_EXEC_IMPL
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::MulOp &op, const ExecutionOptions &opts) {
   auto smallConst = op.getRhs().getDefiningOp<mlir::pphlo::ConstantOp>();
   auto multiplier = op.getLhs();
@@ -272,7 +272,7 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
                                  .getSplatValue<mlir::APFloat>()
                                  .convertToFloat());
       auto eps = kernel::hal::dump_public_as<float>(
-          hctx, kernel::hlo::Epsilon(hctx))[0];
+          sctx, kernel::hlo::Epsilon(sctx))[0];
 
       // Amplify eps to 1/(2^(fxp_bits-2))
       // TODO: Maybe make it configurable?
@@ -287,16 +287,16 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
         auto n = std::pow(2, k);
 
         // n/N
-        auto newRhs = kernel::hlo::Constant(hctx, static_cast<float>(n) / N,
+        auto newRhs = kernel::hlo::Constant(sctx, static_cast<float>(n) / N,
                                             smallConst.getType().getShape());
         // x*n/N
         auto kv = lookupValue(sscope, multiplier, opts);
         // To merge truncation in multiply with next k-bits one, we
         // deliberately pick the ring mul to do a mul *without* truncation
-        auto mulRet = kernel::hal::_mul(hctx, kv, newRhs).setDtype(kv.dtype());
+        auto mulRet = kernel::hal::_mul(sctx, kv, newRhs).setDtype(kv.dtype());
         // truncate(x*n/N, k)
         addValue(sscope, op.getResult(),
-                 kernel::hal::_trunc(hctx, mulRet, hctx->getFxpBits() + k)
+                 kernel::hal::_trunc(sctx, mulRet, sctx->getFxpBits() + k)
                      .setDtype(mulRet.dtype()),
                  opts);
         return;
@@ -305,24 +305,24 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   }
 
   addValue(sscope, op.getResult(),
-           kernel::hlo::Mul(hctx, lookupValue(sscope, op.getLhs(), opts),
+           kernel::hlo::Mul(sctx, lookupValue(sscope, op.getLhs(), opts),
                             lookupValue(sscope, op.getRhs(), opts)),
            opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::DotOp &op, const ExecutionOptions &opts) {
-  auto ret = kernel::hlo::Dot(hctx, lookupValue(sscope, op.getLhs(), opts),
+  auto ret = kernel::hlo::Dot(sctx, lookupValue(sscope, op.getLhs(), opts),
                               lookupValue(sscope, op.getRhs(), opts));
 
   const auto ret_shape =
       op.getResult().getType().dyn_cast<mlir::TensorType>().getShape();
 
-  addValue(sscope, op.getResult(), kernel::hlo::Reshape(hctx, ret, ret_shape),
+  addValue(sscope, op.getResult(), kernel::hlo::Reshape(sctx, ret, ret_shape),
            opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::DotGeneralOp &op, const ExecutionOptions &opts) {
   auto dnum = op.getDotDimensionNumbers();
   // Should in order
@@ -359,25 +359,25 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
     rhs_slice_begin[0] = batch_idx;
     rhs_slice_end[0] = batch_idx + 1;
     auto lhs_slice = kernel::hlo::Reshape(
-        hctx,
-        kernel::hlo::Slice(hctx, lhs, lhs_slice_begin, lhs_slice_end, strides),
+        sctx,
+        kernel::hlo::Slice(sctx, lhs, lhs_slice_begin, lhs_slice_end, strides),
         lhs_slice_shape);
     auto rhs_slice = kernel::hlo::Reshape(
-        hctx,
-        kernel::hlo::Slice(hctx, rhs, rhs_slice_begin, rhs_slice_end, strides),
+        sctx,
+        kernel::hlo::Slice(sctx, rhs, rhs_slice_begin, rhs_slice_end, strides),
         rhs_slice_shape);
     results[batch_idx] = kernel::hlo::Reshape(
-        hctx, kernel::hlo::Dot(hctx, lhs_slice, rhs_slice), ret_slice_shape);
+        sctx, kernel::hlo::Dot(sctx, lhs_slice, rhs_slice), ret_slice_shape);
   }
 
   auto ret_type = op.getResult().getType().dyn_cast<mlir::RankedTensorType>();
   auto ret = kernel::hlo::Reshape(
-      hctx, kernel::hlo::Concatenate(hctx, results, 0), ret_type.getShape());
+      sctx, kernel::hlo::Concatenate(sctx, results, 0), ret_type.getShape());
 
   addValue(sscope, op.getResult(), std::move(ret), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ConvolutionOp &op, const ExecutionOptions &opts) {
   const auto &dnums = op.getDimensionNumbers();
   const size_t num_spatial_dims = dnums.getOutputSpatialDimensions().size();
@@ -416,15 +416,15 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   spu::Value result;
 
   if (dnums.getInputSpatialDimensions().size() == 2) {
-    result = kernel::hlo::Convolution2D(hctx, lhs, rhs, config, ret_shape);
+    result = kernel::hlo::Convolution2D(sctx, lhs, rhs, config, ret_shape);
   } else {
-    result = kernel::hlo::Convolution(hctx, lhs, rhs, config, ret_shape);
+    result = kernel::hlo::Convolution(sctx, lhs, rhs, config, ret_shape);
   }
 
   addValue(sscope, op.getResult(), std::move(result), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::DynamicUpdateSliceOp &op,
              const ExecutionOptions &opts) {
   // Basic idea here, get a ref slice and update the whole slice..
@@ -439,11 +439,11 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
 
   addValue(
       sscope, op.getResult(),
-      kernel::hlo::DynamicUpdateSlice(hctx, operand, update, start_indices),
+      kernel::hlo::DynamicUpdateSlice(sctx, operand, update, start_indices),
       opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::DynamicSliceOp &op, const ExecutionOptions &opts) {
   // Start indices
   auto iter = op.getSliceSizes().getValues<int64_t>();
@@ -456,11 +456,11 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   }
 
   addValue(sscope, op.getResult(),
-           kernel::hlo::DynamicSlice(hctx, operand, slice_size, start_indices),
+           kernel::hlo::DynamicSlice(sctx, operand, slice_size, start_indices),
            opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::GatherOp &op, const ExecutionOptions &opts) {
   // If input is empty, short circuit
   auto operand = lookupValue(sscope, op.getOperand(), opts);
@@ -485,11 +485,11 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
 
   addValue(
       sscope, op.getResult(),
-      kernel::hlo::Gather(hctx, operand, start_indices, config, output_shape),
+      kernel::hlo::Gather(sctx, operand, start_indices, config, output_shape),
       opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::SortOp &op, const ExecutionOptions &opts) {
   auto sort_dim = op.getDimension();
   auto is_stable = op.getIsStable();
@@ -520,10 +520,10 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   }
 
   auto ret = kernel::hlo::Sort(
-      hctx, inputs, sort_dim, is_stable,
+      sctx, inputs, sort_dim, is_stable,
       [&](absl::Span<const spu::Value> inputs) {
         auto ret =
-            runRegion(executor, hctx, sscope, op.getComparator(), inputs);
+            runRegion(executor, sctx, sscope, op.getComparator(), inputs);
         return ret[0];
       },
       spu_return_vis);
@@ -533,7 +533,7 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   }
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::SelectAndScatterOp &op,
              const ExecutionOptions &opts) {
   auto operand = lookupValue(sscope, op.getOperand(), opts);
@@ -565,23 +565,23 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
 
   // auto ret = kernel::hlo::SelectAndScatterNaive(
   auto ret = kernel::hlo::SelectAndScatterExpanded(
-      hctx, operand, source, init_val, window_shape, window_strides,
+      sctx, operand, source, init_val, window_shape, window_strides,
       window_padding,
       [&](const spu::Value &selected, const spu::Value &current) {
-        auto ret = runRegion(executor, hctx, sscope, op.getSelect(),
+        auto ret = runRegion(executor, sctx, sscope, op.getSelect(),
                              {selected, current});
         return ret[0];
       },
       [&](const spu::Value &in, const spu::Value &scatter) {
         auto ret =
-            runRegion(executor, hctx, sscope, op.getScatter(), {in, scatter});
+            runRegion(executor, sctx, sscope, op.getScatter(), {in, scatter});
         return ret[0];
       });
 
   addValue(sscope, op.getResult(), std::move(ret), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::MaxPoolScatterOp &op, const ExecutionOptions &opts) {
   auto scatter_indices = lookupValue(sscope, op.getScatterIndices(), opts);
   auto update = lookupValue(sscope, op.getUpdate(), opts);
@@ -614,22 +614,22 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
       op.getResult().getType().dyn_cast<mlir::RankedTensorType>().getShape();
 
   auto ret =
-      kernel::hlo::MaxPoolScatter(hctx, scatter_indices, update, window_shape,
+      kernel::hlo::MaxPoolScatter(sctx, scatter_indices, update, window_shape,
                                   base_shape, window_strides, window_padding);
 
   addValue(sscope, op.getResult(), std::move(ret), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::CaseOp &op, const ExecutionOptions &opts) {
   std::vector<kernel::hlo::BranchFcnT> branches;
   for (auto &b : op.getBranches()) {
     branches.emplace_back(
-        [&]() { return runRegion(executor, hctx, sscope, b, {}); });
+        [&]() { return runRegion(executor, sctx, sscope, b, {}); });
   }
 
   auto results = kernel::hlo::Case(
-      hctx, lookupValue(sscope, op.getIndex(), opts), branches);
+      sctx, lookupValue(sscope, op.getIndex(), opts), branches);
 
   // Copy output
   for (const auto &ret : llvm::enumerate(op->getResults())) {
@@ -637,17 +637,17 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   }
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::IfOp &op, const ExecutionOptions &opts) {
   auto conditional = lookupValue(sscope, op.getCondition(), opts);
 
   auto results = kernel::hlo::IfElse(
-      hctx, conditional,  //
+      sctx, conditional,  //
       [&]() {
-        return runRegion(executor, hctx, sscope, op.getTrueBranch(), {});
+        return runRegion(executor, sctx, sscope, op.getTrueBranch(), {});
       },
       [&]() {
-        return runRegion(executor, hctx, sscope, op.getFalseBranch(), {});
+        return runRegion(executor, sctx, sscope, op.getFalseBranch(), {});
       });
 
   // Copy output
@@ -656,7 +656,7 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   }
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::WhileOp &op, const ExecutionOptions &opts) {
   // First inputs vectors
   std::vector<spu::Value> inputs;
@@ -668,12 +668,12 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   }
 
   auto ret = kernel::hlo::While(
-      hctx, inputs,  //
+      sctx, inputs,  //
       [&](absl::Span<const spu::Value> inputs) {
-        return runRegion(executor, hctx, sscope, op.getCond(), inputs)[0];
+        return runRegion(executor, sctx, sscope, op.getCond(), inputs)[0];
       },
       [&](absl::Span<const spu::Value> inputs) {
-        return runRegion(executor, hctx, sscope, op.getBody(), inputs);
+        return runRegion(executor, sctx, sscope, op.getBody(), inputs);
       });
 
   for (size_t idx = 0; idx < op->getNumResults(); ++idx) {
@@ -681,7 +681,7 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   }
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::IotaOp &op, const ExecutionOptions &opts) {
   const auto &ret_type =
       op.getOutput().getType().dyn_cast<mlir::RankedTensorType>();
@@ -691,55 +691,55 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   auto ret_el_type = type_tools.getExpressedType(ret_type);
   auto pt_type = getPtTypeFromMlirType(ret_el_type);
 
-  spu::Value iota_ret = kernel::hlo::Iota(hctx, getEncodeType(pt_type), numel);
+  spu::Value iota_ret = kernel::hlo::Iota(sctx, getEncodeType(pt_type), numel);
 
   if (ret_type.getShape().size() > 1) {
     // Need a broadcast
-    iota_ret = kernel::hlo::Broadcast(hctx, iota_ret, ret_type.getShape(), {});
+    iota_ret = kernel::hlo::Broadcast(sctx, iota_ret, ret_type.getShape(), {});
   }
 
   addValue(sscope, op.getOutput(), std::move(iota_ret), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::RemOp &op, const ExecutionOptions &opts) {
   // FIXME: When hal has a remainder, use that
   auto lhs = lookupValue(sscope, op.getLhs(), opts);
   auto rhs = lookupValue(sscope, op.getRhs(), opts);
 
-  auto ret = kernel::hlo::Remainder(hctx, lhs, rhs);
+  auto ret = kernel::hlo::Remainder(sctx, lhs, rhs);
   addValue(sscope, op.getResult(), std::move(ret), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::TransposeOp &op, const ExecutionOptions &opts) {
   addValue(
       sscope, op.getResult(),
-      kernel::hlo::Transpose(hctx, lookupValue(sscope, op.getOperand(), opts),
+      kernel::hlo::Transpose(sctx, lookupValue(sscope, op.getOperand(), opts),
                              convertDenseIntElementAttr(op.getPermutation())),
       opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::BroadcastOp &op, const ExecutionOptions &opts) {
   auto to_shape = op.getType().dyn_cast<mlir::RankedTensorType>().getShape();
   addValue(sscope, op.getResult(),
            kernel::hlo::Broadcast(
-               hctx, lookupValue(sscope, op.getOperand(), opts), to_shape,
+               sctx, lookupValue(sscope, op.getOperand(), opts), to_shape,
                convertDenseIntElementAttr(op.getBroadcastDimensions())),
            opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ReshapeOp &op, const ExecutionOptions &opts) {
   auto to_shape = op.getType().dyn_cast<mlir::RankedTensorType>().getShape();
   addValue(sscope, op.getResult(),
            kernel::hlo::Reshape(
-               hctx, lookupValue(sscope, op.getOperand(), opts), to_shape),
+               sctx, lookupValue(sscope, op.getOperand(), opts), to_shape),
            opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ConcatenateOp &op, const ExecutionOptions &opts) {
   std::vector<spu::Value> values(op->getNumOperands());
 
@@ -749,20 +749,20 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
 
   // set result
   addValue(sscope, op.getResult(),
-           kernel::hlo::Concatenate(hctx, values, op.getDimension()), opts);
+           kernel::hlo::Concatenate(sctx, values, op.getDimension()), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::SliceOp &op, const ExecutionOptions &opts) {
   addValue(sscope, op.getResult(),
-           kernel::hlo::Slice(hctx, lookupValue(sscope, op.getOperand(), opts),
+           kernel::hlo::Slice(sctx, lookupValue(sscope, op.getOperand(), opts),
                               convertDenseIntElementAttr(op.getStartIndices()),
                               convertDenseIntElementAttr(op.getLimitIndices()),
                               convertDenseIntElementAttr(op.getStrides())),
            opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::PadOp &op, const ExecutionOptions &opts) {
   const auto &operand = lookupValue(sscope, op.getOperand(), opts);
   const size_t operand_rank = operand.shape().size();
@@ -779,21 +779,21 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
                           [](int64_t i) { return i >= 0; }));
 
   addValue(sscope, op.getResult(),
-           kernel::hlo::Pad(hctx, operand, padding_value, edge_padding_low,
+           kernel::hlo::Pad(sctx, operand, padding_value, edge_padding_low,
                             edge_padding_high, interior_padding),
            opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ReverseOp &op, const ExecutionOptions &opts) {
   addValue(
       sscope, op.getResult(),
-      kernel::hlo::Reverse(hctx, lookupValue(sscope, op.getOperand(), opts),
+      kernel::hlo::Reverse(sctx, lookupValue(sscope, op.getOperand(), opts),
                            convertDenseIntElementAttr(op.getDimensions())),
       opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ReduceOp &op, const ExecutionOptions &opts) {
   int64_t num_args = op->getNumOperands() / 2;
   std::vector<int64_t> dimensions_to_reduce =
@@ -811,13 +811,13 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
                    [](int64_t d) { return d == 0; });
 
   std::vector<spu::Value> ret = kernel::hlo::Reduce(
-      hctx, input_args, init_values, dimensions_to_reduce,
+      sctx, input_args, init_values, dimensions_to_reduce,
       [&](absl::Span<const spu::Value> lhs, absl::Span<const spu::Value> rhs) {
         std::vector<spu::Value> operands;
         operands.reserve(lhs.size() + rhs.size());
         operands.insert(operands.end(), lhs.begin(), lhs.end());
         operands.insert(operands.end(), rhs.begin(), rhs.end());
-        return runRegion(executor, hctx, sscope, op.getBody(), operands);
+        return runRegion(executor, sctx, sscope, op.getBody(), operands);
       },
       canIgnoreInitialValue);
 
@@ -825,11 +825,11 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
       op->getResultTypes()[0].dyn_cast<mlir::RankedTensorType>().getShape();
   for (size_t idx = 0; idx < op->getNumResults(); ++idx) {
     addValue(sscope, op->getResult(idx),
-             kernel::hlo::Reshape(hctx, ret[idx], output_shape), opts);
+             kernel::hlo::Reshape(sctx, ret[idx], output_shape), opts);
   }
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ReduceWindowOp &op, const ExecutionOptions &opts) {
   int64_t num_args = op->getNumOperands() / 2;
 
@@ -890,13 +890,13 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   config.base_dilations = base_dilation;
 
   auto rets = kernel::hlo::ReduceWindow(
-      hctx, input_args, init_values, ret_shape, config,
+      sctx, input_args, init_values, ret_shape, config,
       [&](absl::Span<const spu::Value> lhs, absl::Span<const spu::Value> rhs) {
         std::vector<spu::Value> operands;
         operands.reserve(lhs.size() + rhs.size());
         operands.insert(operands.end(), lhs.begin(), lhs.end());
         operands.insert(operands.end(), rhs.begin(), rhs.end());
-        return runRegion(executor, hctx, sscope, op.getBody(), operands);
+        return runRegion(executor, sctx, sscope, op.getBody(), operands);
       },
       std::none_of(window_shape.begin(), window_shape.end(),
                    [](int64_t ws) { return ws == 0; }));
@@ -906,7 +906,7 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   }
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ArgMaxOp &op, const ExecutionOptions &opts) {
   auto window_shape = convertDenseIntElementAttr(op.getWindowDimensions());
 
@@ -957,7 +957,7 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   config.window_padding = window_padding;
   config.base_dilations = base_dilation;
 
-  auto ret = kernel::hlo::ArgMax(hctx, lookupValue(sscope, op.getInput(), opts),
+  auto ret = kernel::hlo::ArgMax(sctx, lookupValue(sscope, op.getInput(), opts),
                                  ret_shape, config);
 
   addValue(sscope, op.getResult(0), ret.first, opts);
@@ -965,7 +965,7 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   addValue(sscope, op.getResult(1), ret.second, opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::SelectOp &op, const ExecutionOptions &opts) {
   auto pred = lookupValue(sscope, op.getPred(), opts);
 
@@ -973,20 +973,20 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   auto on_false = lookupValue(sscope, op.getOnFalse(), opts);
 
   addValue(sscope, op.getResult(),
-           kernel::hlo::Select(hctx, pred, on_true, on_false), opts);
+           kernel::hlo::Select(sctx, pred, on_true, on_false), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::RngOp &op, const ExecutionOptions &opts) {
   auto to_shape = op.getType().dyn_cast<mlir::RankedTensorType>().getShape();
   addValue(
       sscope, op.getResult(),
-      kernel::hlo::Uniform_rand(hctx, lookupValue(sscope, op.getA(), opts),
+      kernel::hlo::Uniform_rand(sctx, lookupValue(sscope, op.getA(), opts),
                                 lookupValue(sscope, op.getB(), opts), to_shape),
       opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ConvertOp &op, const ExecutionOptions &opts) {
   mlir::pphlo::TypeTools tool;
   auto dst_dtype = getDtypeFromMlirType(op.getType());
@@ -996,30 +996,30 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
   auto in = lookupValue(sscope, op.getOperand(), opts);
 
   addValue(sscope, op.getResult(),
-           kernel::hlo::Cast(hctx, in, dst_vtype, dst_dtype), opts);
+           kernel::hlo::Cast(sctx, in, dst_vtype, dst_dtype), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::PreferAOp &op, const ExecutionOptions &opts) {
   auto in = lookupValue(sscope, op.getOperand(), opts);
-  if (hctx->rt_config().protocol() == ProtocolKind::CHEETAH) {
+  if (sctx->config().protocol() == ProtocolKind::CHEETAH) {
     // NOTE(juhou): For 2PC, MulAB uses COT which is efficient and accurate than
     // MulAA that needs HE. Thus we just by-pass the PreferAOp for 2PC.
     addValue(sscope, op.getResult(), in, opts);
     return;
   }
-  auto k0 = kernel::hlo::Cast(hctx, kernel::hlo::Constant(hctx, 0, in.shape()),
+  auto k0 = kernel::hlo::Cast(sctx, kernel::hlo::Constant(sctx, 0, in.shape()),
                               VIS_PUBLIC, in.dtype());
-  addValue(sscope, op.getResult(), kernel::hlo::Add(hctx, in, k0), opts);
+  addValue(sscope, op.getResult(), kernel::hlo::Add(sctx, in, k0), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::SignOp &op, const ExecutionOptions &opts) {
   auto in = lookupValue(sscope, op.getOperand(), opts);
-  addValue(sscope, op.getResult(), kernel::hlo::Sign(hctx, in), opts);
+  addValue(sscope, op.getResult(), kernel::hlo::Sign(sctx, in), opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::BitcastConvertOp &op, const ExecutionOptions &opts) {
   const auto &in_type =
       op.getOperand().getType().dyn_cast<mlir::RankedTensorType>();
@@ -1034,12 +1034,12 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
 
   addValue(
       sscope, op.getResult(),
-      kernel::hlo::Bitcast(hctx, lookupValue(sscope, op.getOperand(), opts),
+      kernel::hlo::Bitcast(sctx, lookupValue(sscope, op.getOperand(), opts),
                            getDtypeFromMlirType(out_type)),
       opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ConstantOp &op, const ExecutionOptions &opts) {
   const auto &val = op.getValue();
   const auto &dea = val.dyn_cast<mlir::DenseElementsAttr>();
@@ -1055,7 +1055,7 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
     if (dea.isSplat()) {
       addValue(
           sscope, op.getResult(),
-          kernel::hlo::Constant(hctx, dea.getSplatValue<bool>(), dst_shape),
+          kernel::hlo::Constant(sctx, dea.getSplatValue<bool>(), dst_shape),
           opts);
     } else {
       std::vector<uint8_t> buf(type.getNumElements());
@@ -1066,7 +1066,7 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
                         dst_shape, makeCompactStrides(dst_shape));
 
       addValue(sscope, op.getResult(),
-               kernel::hlo::Constant(hctx, view, dst_shape), opts);
+               kernel::hlo::Constant(sctx, view, dst_shape), opts);
     }
   } else {
     PtBufferView view(
@@ -1075,35 +1075,35 @@ void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
         dea.isSplat() ? std::vector<int64_t>() : makeCompactStrides(dst_shape));
 
     addValue(sscope, op.getResult(),
-             kernel::hlo::Constant(hctx, view, dst_shape), opts);
+             kernel::hlo::Constant(sctx, view, dst_shape), opts);
   }
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::EpsilonOp &op, const ExecutionOptions &opts) {
-  auto e = kernel::hlo::Epsilon(hctx);
+  auto e = kernel::hlo::Epsilon(sctx);
   auto shape =
       op->getResultTypes()[0].dyn_cast<mlir::RankedTensorType>().getShape();
-  addValue(sscope, op.getResult(), kernel::hlo::Broadcast(hctx, e, shape, {}),
+  addValue(sscope, op.getResult(), kernel::hlo::Broadcast(sctx, e, shape, {}),
            opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::ClampOp &op, const ExecutionOptions &opts) {
   addValue(sscope, op.getResult(),
-           kernel::hlo::Clamp(hctx, lookupValue(sscope, op.getOperand(), opts),
+           kernel::hlo::Clamp(sctx, lookupValue(sscope, op.getOperand(), opts),
                               lookupValue(sscope, op.getMin(), opts),
                               lookupValue(sscope, op.getMax(), opts)),
            opts);
 }
 
-void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope,
+void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::DbgPrintOp &op, const ExecutionOptions &opts) {
-  kernel::hal::dbg_print(hctx, lookupValue(sscope, op.getOperand(), opts));
+  kernel::hal::dbg_print(sctx, lookupValue(sscope, op.getOperand(), opts));
 }
 
 #define DEFINE_UNIMPLEMENTED_OP(OpName)                                     \
-  void execute(OpExecutor *executor, HalContext *hctx, SymbolScope *sscope, \
+  void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope, \
                mlir::pphlo::OpName &, const ExecutionOptions &opts) {       \
     SPU_THROW("Lowered op should not occur at backend");                    \
   }
@@ -1135,21 +1135,21 @@ bool PPHloExecutor::hasKernel(mlir::Operation &op) const {
 }
 
 template <typename OpT, typename... MoreOpT>
-static void dispatchOp(OpExecutor *executor, HalContext *hctx,
+static void dispatchOp(OpExecutor *executor, SPUContext *sctx,
                        SymbolScope *sscope, mlir::Operation &op,
                        const ExecutionOptions &opts) {
   if (auto casted = llvm::dyn_cast<OpT>(op)) {
     // Execute op
     {
       const auto fn_name = op.getName().getStringRef().str();
-      SPU_TRACE_ACTION(GET_TRACER(hctx), (TR_HLO | TR_LAR), ~TR_HLO, fn_name);
-      execute(executor, hctx, sscope, casted, opts);
+      SPU_TRACE_ACTION(GET_TRACER(sctx), (TR_HLO | TR_LAR), ~TR_HLO, fn_name);
+      execute(executor, sctx, sscope, casted, opts);
     }
 
     // currently we only support config verifier statically.
     constexpr bool kEnableXlaVerifier = false;
     if (kEnableXlaVerifier) {
-      PPHloVerifier verifier(hctx);
+      PPHloVerifier verifier(sctx);
       // handle mixed (int, fxp) multiplication
       if constexpr (std::is_same_v<OpT, mlir::pphlo::MulOp> or
                     std::is_same_v<OpT, mlir::pphlo::DotOp> or
@@ -1164,10 +1164,10 @@ static void dispatchOp(OpExecutor *executor, HalContext *hctx,
             type_tool.getExpressedType(casted.getResult().getType());
 
         if (lhs_type != ret_type) {
-          lhs = kernel::hlo::Cast(hctx, lhs, lhs.vtype(), ret.dtype());
+          lhs = kernel::hlo::Cast(sctx, lhs, lhs.vtype(), ret.dtype());
         }
         if (rhs_type != ret_type) {
-          rhs = kernel::hlo::Cast(hctx, rhs, rhs.vtype(), ret.dtype());
+          rhs = kernel::hlo::Cast(sctx, rhs, rhs.vtype(), ret.dtype());
         }
 
         verifier.verify(casted, {lhs, rhs}, {ret});
@@ -1190,12 +1190,12 @@ static void dispatchOp(OpExecutor *executor, HalContext *hctx,
       SPU_THROW("Unhandled mlir op {} at {}", mlirObjectToString(op),
                 mlirObjectToString(op.getLoc()));
     } else {
-      dispatchOp<MoreOpT...>(executor, hctx, sscope, op, opts);
+      dispatchOp<MoreOpT...>(executor, sctx, sscope, op, opts);
     }
   }
 }
 
-void PPHloExecutor::runKernelImpl(HalContext *hctx, SymbolScope *sscope,
+void PPHloExecutor::runKernelImpl(SPUContext *sctx, SymbolScope *sscope,
                                   mlir::Operation &op,
                                   const ExecutionOptions &opts) {
   if (opts.do_log_execution) {
@@ -1204,7 +1204,7 @@ void PPHloExecutor::runKernelImpl(HalContext *hctx, SymbolScope *sscope,
   dispatchOp<
 #define GET_OP_LIST
 #include "libspu/dialect/pphlo_ops.cc.inc"
-      >(this, hctx, sscope, op, opts);
+      >(this, sctx, sscope, op, opts);
 }
 
 void PPHloExecutor::checkType(mlir::Type mlir_type, const spu::Value &v) const {
diff --git a/libspu/device/pphlo/pphlo_executor.h b/libspu/device/pphlo/pphlo_executor.h
index 2bfc48e1..2e1284ad 100644
--- a/libspu/device/pphlo/pphlo_executor.h
+++ b/libspu/device/pphlo/pphlo_executor.h
@@ -16,11 +16,11 @@
 
 #include <memory>
 
+#include "libspu/core/context.h"
 #include "libspu/device/executor.h"
 #include "libspu/device/pphlo/pphlo_verifier.h"
 #include "libspu/dialect/pphlo_ops.h"
 #include "libspu/dialect/pphlo_types.h"
-#include "libspu/kernel/context.h"
 #include "libspu/kernel/hlo/casting.h"
 
 namespace spu::device::pphlo {
@@ -33,7 +33,7 @@ class PPHloExecutor : public OpExecutor {
   bool hasKernel(mlir::Operation &op) const override;
 
   // run a kernel in a given region.
-  void runKernelImpl(HalContext *hctx, SymbolScope *sscope, mlir::Operation &op,
+  void runKernelImpl(SPUContext *sctx, SymbolScope *sscope, mlir::Operation &op,
                      const ExecutionOptions &opts) override;
 };
 
diff --git a/libspu/device/pphlo/pphlo_verifier.cc b/libspu/device/pphlo/pphlo_verifier.cc
index bd125d10..a0af24bb 100644
--- a/libspu/device/pphlo/pphlo_verifier.cc
+++ b/libspu/device/pphlo/pphlo_verifier.cc
@@ -62,7 +62,7 @@ mlir::TensorType buildMLIRType(mlir::MLIRContext *mlir_ctx,
 }
 
 mlir::stablehlo::Tensor convertToStablehloTensor(mlir::MLIRContext *mlir_ctx,
-                                                 HalContext *ctx,
+                                                 SPUContext *ctx,
                                                  const spu::Value &v) {
   NdArrayRef arr;
   if (v.isSecret()) {
@@ -192,7 +192,7 @@ std::string mlirTypeToString(mlir::Type type) {
   return os.str();
 }
 
-bool verifyEqual(HalContext *ctx, const mlir::stablehlo::Tensor &xla_ret,
+bool verifyEqual(SPUContext *ctx, const mlir::stablehlo::Tensor &xla_ret,
                  const mlir::stablehlo::Tensor &spu_ret) {
   if (xla_ret.getType() != spu_ret.getType()) {
     SPDLOG_INFO("Answer has a type mismatch, xla type = {}, spu type = {}",
@@ -238,7 +238,7 @@ bool verifyEqual(HalContext *ctx, const mlir::stablehlo::Tensor &xla_ret,
 
 }  // namespace
 
-PPHloVerifier::PPHloVerifier(HalContext *ctx) : ctx_(ctx) {
+PPHloVerifier::PPHloVerifier(SPUContext *ctx) : ctx_(ctx) {
   mlir_ctx_.getOrLoadDialect<mlir::stablehlo::StablehloDialect>();
 }
 
diff --git a/libspu/device/pphlo/pphlo_verifier.h b/libspu/device/pphlo/pphlo_verifier.h
index 9190e98a..755aa81a 100644
--- a/libspu/device/pphlo/pphlo_verifier.h
+++ b/libspu/device/pphlo/pphlo_verifier.h
@@ -14,22 +14,22 @@
 
 #pragma once
 
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 #include "libspu/dialect/pphlo_dialect.h"
 #include "libspu/dialect/pphlo_ops.h"
 #include "libspu/dialect/pphlo_types.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
 
 namespace spu::device::pphlo {
 
 class PPHloVerifier {
  private:
-  HalContext *ctx_{nullptr};
+  SPUContext *ctx_{nullptr};
   mlir::MLIRContext mlir_ctx_;
   std::function<void(bool)> mismatch_handler_{[](bool) {}};
 
  public:
-  explicit PPHloVerifier(HalContext *ctx);
+  explicit PPHloVerifier(SPUContext *ctx);
 
   void setMismatchHandler(std::function<void(bool)> f) {
     mismatch_handler_ = std::move(f);
diff --git a/libspu/device/symbol_table.h b/libspu/device/symbol_table.h
index 39d4207d..2619456c 100644
--- a/libspu/device/symbol_table.h
+++ b/libspu/device/symbol_table.h
@@ -17,7 +17,7 @@
 #include <string>
 #include <unordered_map>
 
-#include "libspu/kernel/value.h"
+#include "libspu/core/value.h"
 
 #include "libspu/device/device.pb.h"
 
diff --git a/libspu/kernel/BUILD.bazel b/libspu/kernel/BUILD.bazel
index 6069aa00..c65da3a0 100644
--- a/libspu/kernel/BUILD.bazel
+++ b/libspu/kernel/BUILD.bazel
@@ -17,33 +17,17 @@ load("//bazel:spu.bzl", "spu_cc_library", "spu_cc_test")
 package(default_visibility = ["//visibility:public"])
 
 spu_cc_library(
-    name = "value",
-    srcs = ["value.cc"],
-    hdrs = ["value.h"],
+    name = "test_util",
+    testonly = True,
+    srcs = ["test_util.cc"],
+    hdrs = ["test_util.h"],
     deps = [
-        "//libspu/core:ndarray_ref",
-        "//libspu/core:prelude",
-        "//libspu/core:shape_util",
-    ],
-)
-
-spu_cc_test(
-    name = "value_test",
-    srcs = ["value_test.cc"],
-    deps = [
-        "//libspu/kernel:value",
-    ],
-)
-
-spu_cc_library(
-    name = "context",
-    srcs = ["context.cc"],
-    hdrs = ["context.h"],
-    deps = [
-        "//libspu/core",
-        "//libspu/core:trace",
-        "//libspu/kernel:value",  # FIXME: each module depends on value
+        "//libspu/core:context",
+        "//libspu/core:value",
+        "//libspu/kernel/hal:constants",  # BAD
+        "//libspu/kernel/hal:prot_wrapper",  # BAD
+        "//libspu/kernel/hal:public_helper",  # BAD
         "//libspu/mpc:factory",
-        "@yacl//yacl/link",
+        "@com_google_googletest//:gtest",
     ],
 )
diff --git a/libspu/kernel/context.cc b/libspu/kernel/context.cc
deleted file mode 100644
index c10998fd..00000000
--- a/libspu/kernel/context.cc
+++ /dev/null
@@ -1,40 +0,0 @@
-// Copyright 2021 Ant Group Co., Ltd.
-//
-// 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 "libspu/kernel/context.h"
-
-#include "libspu/core/config.h"
-#include "libspu/mpc/factory.h"
-
-namespace spu {
-
-HalContext::HalContext(const RuntimeConfig& config,
-                       const std::shared_ptr<yacl::link::Context>& lctx)
-    : rt_config_(makeFullRuntimeConfig(config)),
-      lctx_(lctx),
-      prot_(mpc::Factory::CreateCompute(rt_config_, lctx)) {}
-
-std::unique_ptr<HalContext> HalContext::fork() {
-  auto new_hctx = std::unique_ptr<HalContext>(new HalContext);
-
-  new_hctx->rt_config_ = rt_config_;
-  if (lctx_) {
-    new_hctx->lctx_ = lctx_->Spawn();
-  }
-  new_hctx->prot_ = prot_->fork();
-
-  return new_hctx;
-}
-
-}  // namespace spu
diff --git a/libspu/kernel/context.h b/libspu/kernel/context.h
deleted file mode 100644
index 05f3a328..00000000
--- a/libspu/kernel/context.h
+++ /dev/null
@@ -1,74 +0,0 @@
-// Copyright 2021 Ant Group Co., Ltd.
-//
-// 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.
-
-#pragma once
-
-#include <memory>
-#include <random>
-
-#include "yacl/link/link.h"
-
-#include "libspu/core/prelude.h"
-#include "libspu/core/trace.h"
-#include "libspu/mpc/object.h"
-
-#include "libspu/spu.pb.h"
-
-namespace spu {
-
-// The hal evaluation context for all spu operators.
-class HalContext final {
-  RuntimeConfig rt_config_;
-
-  std::shared_ptr<yacl::link::Context> lctx_;
-
-  std::unique_ptr<mpc::Object> prot_;
-
-  HalContext() = default;
-
- public:
-  explicit HalContext(const RuntimeConfig& config,
-                      const std::shared_ptr<yacl::link::Context>& lctx);
-
-  HalContext(const HalContext& other) = delete;
-  HalContext& operator=(const HalContext& other) = delete;
-
-  // all parties get a 'corresponding' hal context when forked.
-  std::unique_ptr<HalContext> fork();
-
-  const std::string& id() { return prot_->id(); }
-  const std::string& pid() { return prot_->pid(); }
-
-  HalContext(HalContext&& other) = default;
-
-  //
-  const std::shared_ptr<yacl::link::Context>& lctx() const { return lctx_; }
-
-  mpc::Object* prot() const { return prot_.get(); }
-
-  // Return current working fixed point fractional bits.
-  size_t getFxpBits() const {
-    auto fbits = rt_config_.fxp_fraction_bits();
-    SPU_ENFORCE(fbits != 0);
-    return fbits;
-  }
-
-  // Return current working field of MPC engine.
-  FieldType getField() const { return rt_config_.field(); }
-
-  // Return current working runtime config.
-  const RuntimeConfig& rt_config() const { return rt_config_; }
-};
-
-}  // namespace spu
diff --git a/libspu/kernel/hal/BUILD.bazel b/libspu/kernel/hal/BUILD.bazel
index 022d5504..630b2d28 100644
--- a/libspu/kernel/hal/BUILD.bazel
+++ b/libspu/kernel/hal/BUILD.bazel
@@ -29,7 +29,7 @@ spu_cc_library(
         ":random",
         ":shape_ops",
         ":type_cast",
-        "//libspu/kernel:value",
+        "//libspu/core:value",
     ],
 )
 
@@ -38,7 +38,9 @@ spu_cc_library(
     srcs = ["prot_wrapper.cc"],
     hdrs = ["prot_wrapper.h"],
     deps = [
-        "//libspu/kernel:context",
+        "//libspu/core:context",
+        "//libspu/core:value",
+        "//libspu/mpc:api",
     ],
 )
 
@@ -48,7 +50,7 @@ spu_cc_library(
     hdrs = ["ring.h"],
     deps = [
         ":prot_wrapper",
-        "//libspu/kernel:context",
+        "//libspu/core:context",
     ],
 )
 
@@ -68,7 +70,7 @@ spu_cc_test(
     deps = [
         ":constants",
         ":integer",
-        ":test_util",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -78,7 +80,7 @@ spu_cc_test(
     deps = [
         ":constants",
         ":ring",
-        ":test_util",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -87,7 +89,9 @@ spu_cc_library(
     srcs = ["fxp_cleartext.cc"],
     hdrs = ["fxp_cleartext.h"],
     deps = [
-        "//libspu/kernel:context",
+        "//libspu/core:context",
+        "//libspu/core:encoding",
+        "//libspu/core:value",
     ],
 )
 
@@ -106,7 +110,7 @@ spu_cc_test(
     srcs = ["fxp_base_test.cc"],
     deps = [
         ":fxp_base",
-        ":test_util",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -126,7 +130,7 @@ spu_cc_test(
     srcs = ["fxp_approx_test.cc"],
     deps = [
         ":fxp_approx",
-        ":test_util",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -146,8 +150,11 @@ spu_cc_library(
     deps = [
         ":prot_wrapper",
         ":ring",
+        "//libspu/core:context",
+        "//libspu/core:encoding",
         "//libspu/core:pt_buffer_view",
-        "//libspu/kernel:context",
+        "//libspu/core:value",
+        "//libspu/mpc/common:pub2k",  # TODO: this is a bad reference
     ],
 )
 
@@ -156,7 +163,7 @@ spu_cc_test(
     srcs = ["constants_test.cc"],
     deps = [
         ":constants",
-        ":test_util",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -165,25 +172,13 @@ spu_cc_library(
     srcs = ["public_helper.cc"],
     hdrs = ["public_helper.h"],
     deps = [
+        "//libspu/core:context",
         "//libspu/core:encoding",
-        "//libspu/kernel:context",
+        "//libspu/core:value",
         "//libspu/mpc/common:pub2k",  # TODO: this is a bad reference
     ],
 )
 
-spu_cc_library(
-    name = "test_util",
-    testonly = True,
-    srcs = ["test_util.cc"],
-    hdrs = ["test_util.h"],
-    deps = [
-        ":constants",
-        ":public_helper",
-        "//libspu/kernel:context",
-        "@com_google_googletest//:gtest",
-    ],
-)
-
 spu_cc_library(
     name = "debug",
     srcs = ["debug.cc"],
@@ -215,8 +210,8 @@ spu_cc_test(
     srcs = ["polymorphic_test.cc"],
     deps = [
         ":polymorphic",
-        ":test_util",
-        "//libspu/kernel:context",
+        "//libspu/core:context",
+        "//libspu/kernel:test_util",
         "//libspu/mpc/utils:linalg",
     ],
 )
@@ -226,8 +221,8 @@ spu_cc_test(
     srcs = ["array_element_wise_test.cc"],
     deps = [
         ":polymorphic",
-        ":test_util",
-        "//libspu/kernel:context",
+        "//libspu/core:context",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -246,8 +241,8 @@ spu_cc_test(
     name = "type_cast_test",
     srcs = ["type_cast_test.cc"],
     deps = [
-        ":test_util",
         ":type_cast",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -258,8 +253,8 @@ spu_cc_library(
     deps = [
         ":constants",
         ":public_helper",
+        "//libspu/core:context",
         "//libspu/core:prelude",
-        "//libspu/kernel:context",
     ],
 )
 
@@ -270,9 +265,9 @@ spu_cc_library(
     deps = [
         ":ring",
         ":type_cast",
+        "//libspu/core:context",
         "//libspu/core:prelude",
         "//libspu/core:vectorize",
-        "//libspu/kernel:context",
     ],
 )
 
@@ -281,6 +276,6 @@ spu_cc_test(
     srcs = ["shape_ops_test.cc"],
     deps = [
         ":shape_ops",
-        ":test_util",
+        "//libspu/kernel:test_util",
     ],
 )
diff --git a/libspu/kernel/hal/constants.cc b/libspu/kernel/hal/constants.cc
index b934dbf6..f62d46d1 100644
--- a/libspu/kernel/hal/constants.cc
+++ b/libspu/kernel/hal/constants.cc
@@ -29,7 +29,7 @@ namespace {
 // make a public typed value.
 //
 // FIXME: this is a abstraction leakage, we should NOT invoke Pub2kTy directly.
-Value make_pub2k(HalContext* ctx, const PtBufferView& bv) {
+Value make_pub2k(SPUContext* ctx, const PtBufferView& bv) {
   SPU_TRACE_HAL_DISP(ctx, bv);
 
   NdArrayRef raw = convertToNdArray(bv);
@@ -69,7 +69,7 @@ bool canImplicitCastTo(DataType frm, DataType to) {
 
 }  // namespace
 
-Value constant(HalContext* ctx, PtBufferView init, DataType dtype,
+Value constant(SPUContext* ctx, PtBufferView init, DataType dtype,
                ShapeView shape) {
   SPU_TRACE_HAL_DISP(ctx, init, dtype, shape);
 
@@ -102,7 +102,7 @@ Value constant(HalContext* ctx, PtBufferView init, DataType dtype,
   return Value(result.data().broadcast_to(shape, {}), result.dtype());
 }
 
-spu::Value zeros(HalContext* ctx, DataType dtype,
+spu::Value zeros(SPUContext* ctx, DataType dtype,
                  absl::Span<const int64_t> shape) {
   if (dtype == DT_FXP) {
     return constant(ctx, 0.0, DT_FXP, shape);
@@ -111,7 +111,7 @@ spu::Value zeros(HalContext* ctx, DataType dtype,
   }
 }
 
-Value iota(HalContext* ctx, DataType dtype, int64_t numel) {
+Value iota(SPUContext* ctx, DataType dtype, int64_t numel) {
   return DISPATCH_ALL_NONE_BOOL_PT_TYPES(getDecodeType(dtype), "iota", [&]() {
     std::vector<ScalarT> arr(numel);
     std::iota(arr.begin(), arr.end(), 0);
@@ -119,7 +119,7 @@ Value iota(HalContext* ctx, DataType dtype, int64_t numel) {
   });
 }
 
-Value epsilon(HalContext* ctx, absl::Span<const int64_t> shape) {
+Value epsilon(SPUContext* ctx, absl::Span<const int64_t> shape) {
   return _constant(ctx, static_cast<int128_t>(1), shape).asFxp();
 }
 
diff --git a/libspu/kernel/hal/constants.h b/libspu/kernel/hal/constants.h
index c6d6c91c..9e202210 100644
--- a/libspu/kernel/hal/constants.h
+++ b/libspu/kernel/hal/constants.h
@@ -14,35 +14,33 @@
 
 #pragma once
 
+#include "libspu/core/context.h"
 #include "libspu/core/pt_buffer_view.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/shape_util.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hal {
 
-// TODO: move this to shape_utils.
-using ShapeView = absl::Span<int64_t const>;
-
 // Returns a SPU value from given C/C++ buffer.
 //
 // The result visibility is public.
-Value constant(HalContext* ctx, PtBufferView init, DataType dtype,
+Value constant(SPUContext* ctx, PtBufferView init, DataType dtype,
                ShapeView shape = {});
 
 // Returns a SPU zero value, which is equal to
 //  constant(ctx, 0, dtype, shape);
 //
 // The result visibility is public.
-Value zeros(HalContext* ctx, DataType dtype, ShapeView shape = {});
+Value zeros(SPUContext* ctx, DataType dtype, ShapeView shape = {});
 
 // Returns a one-dimentional value.
 //
 // The result visibility is public.
-Value iota(HalContext* ctx, DataType dtype, int64_t numel);
+Value iota(SPUContext* ctx, DataType dtype, int64_t numel);
 
 // Returns the SPU epsilon, the positive distance between two fixed point value.
 //
 // The result visibility is public.
-Value epsilon(HalContext* ctx, absl::Span<const int64_t> shape = {});
+Value epsilon(SPUContext* ctx, absl::Span<const int64_t> shape = {});
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/debug.cc b/libspu/kernel/hal/debug.cc
index 4190d00d..f2b1825e 100644
--- a/libspu/kernel/hal/debug.cc
+++ b/libspu/kernel/hal/debug.cc
@@ -21,7 +21,7 @@
 
 namespace spu::kernel::hal {
 
-void dbg_print(HalContext* ctx, const Value& v) {
+void dbg_print(SPUContext* ctx, const Value& v) {
   if (v.isPublic()) {
     std::stringstream ss;
     if (v.isFxp()) {
diff --git a/libspu/kernel/hal/debug.h b/libspu/kernel/hal/debug.h
index 9d251d1f..840daf85 100644
--- a/libspu/kernel/hal/debug.h
+++ b/libspu/kernel/hal/debug.h
@@ -14,13 +14,13 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hal {
 
 /// Debug print a value, this action may reveal secret.
 // @param v, value to print
-void dbg_print(HalContext* ctx, const Value& v);
+void dbg_print(SPUContext* ctx, const Value& v);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/fxp_approx.cc b/libspu/kernel/hal/fxp_approx.cc
index 606a7b9a..05269744 100644
--- a/libspu/kernel/hal/fxp_approx.cc
+++ b/libspu/kernel/hal/fxp_approx.cc
@@ -28,7 +28,7 @@ namespace spu::kernel::hal {
 namespace {
 
 // simple convenient function.
-Value f_constant(HalContext* ctx, const PtBufferView& init,
+Value f_constant(SPUContext* ctx, const PtBufferView& init,
                  absl::Span<int64_t const> shape) {
   return constant(ctx, init, DT_FXP, shape);
 }
@@ -49,7 +49,7 @@ namespace detail {
 //          + x^3 * 0.1 *10
 // log2(x) = p2524(x) / q2524(x)
 //
-Value log2_pade_approx_for_normalized(HalContext* ctx, const Value& x) {
+Value log2_pade_approx_for_normalized(SPUContext* ctx, const Value& x) {
   const auto x2 = f_square(ctx, x);
   const auto x3 = f_mul(ctx, x2, x);
 
@@ -80,10 +80,10 @@ Value log2_pade_approx_for_normalized(HalContext* ctx, const Value& x) {
 // Chapter 5 Exponentiation and Logarithms
 // Benchmarking Privacy Preserving Scientific Operations
 // https://www.esat.kuleuven.be/cosic/publications/article-3013.pdf
-Value log2_pade_approx(HalContext* ctx, const Value& x) {
+Value log2_pade_approx(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
-  const size_t bit_width = SizeOf(ctx->rt_config().field()) * 8;
+  const size_t bit_width = SizeOf(ctx->config().field()) * 8;
   auto k = _popcount(ctx, _prefix_or(ctx, x), bit_width);
 
   const size_t num_fxp_bits = ctx->getFxpBits();
@@ -110,7 +110,7 @@ Value log2_pade_approx(HalContext* ctx, const Value& x) {
 // Approximates the natural logarithm using 8th order modified
 // Householder iterations. This approximation is accurate within 2% relative
 // error on [0.0001, 250].
-Value log_householder_approx(HalContext* ctx, const Value& x) {
+Value log_householder_approx(SPUContext* ctx, const Value& x) {
   Value term_1 = f_div(ctx, x, f_constant(ctx, 120.0, x.shape()));
   Value term_2 = f_mul(
       ctx,
@@ -123,14 +123,14 @@ Value log_householder_approx(HalContext* ctx, const Value& x) {
       f_add(ctx, f_sub(ctx, term_1, term_2), f_constant(ctx, 3.0, x.shape()));
 
   std::vector<Value> coeffs;
-  const size_t fxp_log_orders = ctx->rt_config().fxp_log_orders();
+  const size_t fxp_log_orders = ctx->config().fxp_log_orders();
   SPU_ENFORCE(fxp_log_orders != 0, "fxp_log_orders should not be {}",
               fxp_log_orders);
   for (size_t i = 0; i < fxp_log_orders; i++) {
     coeffs.emplace_back(f_constant(ctx, 1.0 / (1.0 + i), x.shape()));
   }
 
-  const size_t num_iters = ctx->rt_config().fxp_log_iters();
+  const size_t num_iters = ctx->config().fxp_log_iters();
   SPU_ENFORCE(num_iters != 0, "fxp_log_iters should not be {}", num_iters);
   for (size_t i = 0; i < num_iters; i++) {
     Value h = f_sub(ctx, f_constant(ctx, 1.0, x.shape()),
@@ -143,8 +143,8 @@ Value log_householder_approx(HalContext* ctx, const Value& x) {
 
 // see https://lvdmaaten.github.io/publications/papers/crypten.pdf
 //   exp(x) = (1 + x / n) ^ n, when n is infinite large.
-Value exp_taylor_series(HalContext* ctx, const Value& x) {
-  const size_t fxp_exp_iters = ctx->rt_config().fxp_exp_iters();
+Value exp_taylor_series(SPUContext* ctx, const Value& x) {
+  const size_t fxp_exp_iters = ctx->config().fxp_exp_iters();
   SPU_ENFORCE(fxp_exp_iters != 0, "fxp_exp_iters should not be {}",
               fxp_exp_iters);
 
@@ -165,7 +165,7 @@ Value exp_taylor_series(HalContext* ctx, const Value& x) {
 //             + x^3 * 0.555040686204663 / 10
 //             + x^4 * 0.961834122588046 / 100
 //             + x^5 * 0.133273035928143 / 100
-Value exp2_pade_approx_for_positive_pure_decimal(HalContext* ctx,
+Value exp2_pade_approx_for_positive_pure_decimal(SPUContext* ctx,
                                                  const Value& x) {
   auto x2 = f_mul(ctx, x, x);
   auto x3 = f_mul(ctx, x, x2);
@@ -194,7 +194,7 @@ Value exp2_pade_approx_for_positive_pure_decimal(HalContext* ctx,
 // https://www.esat.kuleuven.be/cosic/publications/article-3013.pdf
 // NOTE(junfeng): The valid integer bits of x is 5. Otherwise, the output is
 // incorrect.
-Value exp2_pade_approx(HalContext* ctx, const Value& x) {
+Value exp2_pade_approx(SPUContext* ctx, const Value& x) {
   const size_t fbits = ctx->getFxpBits();
   const auto k1 = _constant(ctx, 1U, x.shape());
   // TODO(junfeng): Make int_bits configurable.
@@ -232,7 +232,7 @@ Value exp2_pade_approx(HalContext* ctx, const Value& x) {
                _mul(ctx, x_msb, f_sub(ctx, ret_reciprocal, ret)).asFxp());
 }
 
-Value exp_pade_approx(HalContext* ctx, const Value& x) {
+Value exp_pade_approx(SPUContext* ctx, const Value& x) {
   return f_exp2(
       ctx, f_mul(ctx, x, f_constant(ctx, std::log2(std::exp(1)), x.shape())));
 }
@@ -241,7 +241,7 @@ Value exp_pade_approx(HalContext* ctx, const Value& x) {
 // https://www.wolframalpha.com/input?i=Pade+approximation+tanh%28x%29+order+5%2C5.
 // tanh(x) = (x + x^3 / 9.0 + x^5 /945.0) /
 //           (1 + 4 * x^2 / 9.0 + x^4 / 63.0)
-Value tanh_pade_approx(HalContext* ctx, const Value& x) {
+Value tanh_pade_approx(SPUContext* ctx, const Value& x) {
   const auto x_2 = f_square(ctx, x);
   const auto x_3 = f_mul(ctx, x_2, x);
   const auto x_4 = f_square(ctx, x_2);
@@ -262,7 +262,7 @@ Value tanh_pade_approx(HalContext* ctx, const Value& x) {
 
 }  // namespace detail
 
-Value f_exp(HalContext* ctx, const Value& x) {
+Value f_exp(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
@@ -271,7 +271,7 @@ Value f_exp(HalContext* ctx, const Value& x) {
     return f_exp_p(ctx, x);
   }
 
-  switch (ctx->rt_config().fxp_exp_mode()) {
+  switch (ctx->config().fxp_exp_mode()) {
     case RuntimeConfig::EXP_DEFAULT:
     case RuntimeConfig::EXP_TAYLOR:
       return detail::exp_taylor_series(ctx, x);
@@ -287,11 +287,11 @@ Value f_exp(HalContext* ctx, const Value& x) {
     }
     default:
       SPU_THROW("unexpected exp approxmation method {}",
-                ctx->rt_config().fxp_exp_mode());
+                ctx->config().fxp_exp_mode());
   }
 }
 
-Value f_log(HalContext* ctx, const Value& x) {
+Value f_log(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
@@ -300,7 +300,7 @@ Value f_log(HalContext* ctx, const Value& x) {
     return f_log_p(ctx, x);
   }
 
-  switch (ctx->rt_config().fxp_log_mode()) {
+  switch (ctx->config().fxp_log_mode()) {
     case RuntimeConfig::LOG_DEFAULT:
     case RuntimeConfig::LOG_PADE:
       return f_mul(ctx, f_constant(ctx, std::log(2.0), x.shape()),
@@ -309,11 +309,11 @@ Value f_log(HalContext* ctx, const Value& x) {
       return detail::log_householder_approx(ctx, x);
     default:
       SPU_THROW("unlogected log approxmation method {}",
-                ctx->rt_config().fxp_log_mode());
+                ctx->config().fxp_log_mode());
   }
 }
 
-Value f_log1p(HalContext* ctx, const Value& x) {
+Value f_log1p(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
@@ -321,7 +321,7 @@ Value f_log1p(HalContext* ctx, const Value& x) {
   return f_log(ctx, f_add(ctx, f_constant(ctx, 1.0, x.shape()), x));
 }
 
-Value f_log2(HalContext* ctx, const Value& x) {
+Value f_log2(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
@@ -329,19 +329,19 @@ Value f_log2(HalContext* ctx, const Value& x) {
   return detail::log2_pade_approx(ctx, x).asFxp();
 }
 
-Value f_exp2(HalContext* ctx, const Value& x) {
+Value f_exp2(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   return detail::exp2_pade_approx(ctx, x);
 }
 
-Value f_tanh(HalContext* ctx, const Value& x) {
+Value f_tanh(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   return detail::tanh_pade_approx(ctx, x);
 }
 
-static Value rsqrt_init_guess(HalContext* ctx, const Value& x, const Value& z) {
+static Value rsqrt_init_guess(SPUContext* ctx, const Value& x, const Value& z) {
   SPU_TRACE_HAL_LEAF(ctx, x, z);
   const size_t f = ctx->getFxpBits();
 
@@ -354,7 +354,7 @@ static Value rsqrt_init_guess(HalContext* ctx, const Value& x, const Value& z) {
   // let rsqrt(u) = 26.02942339 * u^4 - 49.86605845 * u^3 + 38.4714796 * u^2
   // - 15.47994394 * u + 4.14285016
   spu::Value r;
-  if (!ctx->rt_config().enable_lower_accuracy_rsqrt()) {
+  if (!ctx->config().enable_lower_accuracy_rsqrt()) {
     std::vector<Value> coeffs = {f_constant(ctx, -15.47994394, x.shape()),
                                  f_constant(ctx, 38.4714796, x.shape()),
                                  f_constant(ctx, -49.86605845, x.shape()),
@@ -371,7 +371,7 @@ static Value rsqrt_init_guess(HalContext* ctx, const Value& x, const Value& z) {
   return r;
 }
 
-static Value rsqrt_comp(HalContext* ctx, const Value& x, const Value& z) {
+static Value rsqrt_comp(SPUContext* ctx, const Value& x, const Value& z) {
   SPU_TRACE_HAL_LEAF(ctx, x, z);
 
   const size_t k = SizeOf(ctx->getField()) * 8;
@@ -422,7 +422,7 @@ static Value rsqrt_comp(HalContext* ctx, const Value& x, const Value& z) {
   return _mul(ctx, _mux(ctx, b, c0, c1), a_rev);
 }
 
-static Value rsqrt_np2(HalContext* ctx, const Value& x) {
+static Value rsqrt_np2(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   // let e = NP2(x), z = 2^(e+f)
@@ -436,7 +436,7 @@ static Value rsqrt_np2(HalContext* ctx, const Value& x) {
 //  2. get a nice approximation for u part.
 //  3. get the compensation for 2^(e + 1) part.
 //  4. multiple two parts and get the result.
-Value f_rsqrt(HalContext* ctx, const Value& x) {
+Value f_rsqrt(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   // let e = NP2(x) , z = 2^(e+f)
@@ -444,9 +444,10 @@ Value f_rsqrt(HalContext* ctx, const Value& x) {
 
   // TODO: we should avoid fork context in hal layer, it will make global
   // scheduling harder and also make profiling harder.
-  if (ctx->rt_config().experimental_enable_intra_op_par()) {
+  if (ctx->config().experimental_enable_intra_op_par()) {
     auto sub_ctx = ctx->fork();
-    auto r = std::async(rsqrt_init_guess, sub_ctx.get(), x, z);
+    auto r = std::async(rsqrt_init_guess,
+                        dynamic_cast<SPUContext*>(sub_ctx.get()), x, z);
     auto comp = rsqrt_comp(ctx, x, z);
     return _trunc(ctx, _mul(ctx, r.get(), comp)).asFxp();
   } else {
@@ -463,7 +464,7 @@ Value f_rsqrt(HalContext* ctx, const Value& x) {
 // Goldschmidt iteration, needs an initial approximation of sqrt_inv(x).
 // In the end, g is an approximation of sqrt(x) while h is an approximation of
 // 1 / (2 * sqrt(x)).
-Value f_sqrt(HalContext* ctx, const Value& x) {
+Value f_sqrt(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   const auto c0 = f_constant(ctx, 0.5, x.shape());
@@ -487,20 +488,20 @@ Value f_sqrt(HalContext* ctx, const Value& x) {
 
 namespace {
 
-Value sigmiod_real(HalContext* ctx, const Value& x) {
+Value sigmiod_real(SPUContext* ctx, const Value& x) {
   // f(x) = 1/(1+exp(-x))
   const auto c1 = f_constant(ctx, 1.0F, x.shape());
   return f_reciprocal(ctx, f_add(ctx, c1, f_exp(ctx, f_negate(ctx, x))));
 }
 
-Value sigmiod_mm1(HalContext* ctx, const Value& x) {
+Value sigmiod_mm1(SPUContext* ctx, const Value& x) {
   // SigmoidMM1: f(x) = 0.5 + 0.125 * x
   const auto c1 = f_constant(ctx, 0.5, x.shape());
   const auto c2 = f_constant(ctx, 0.125, x.shape());
   return f_add(ctx, c1, f_mul(ctx, c2, x));
 }
 
-Value sigmiod_seg3(HalContext* ctx, const Value& x) {
+Value sigmiod_seg3(SPUContext* ctx, const Value& x) {
   // f(x) = 0.5 + 0.125x if -4 <= x <= 4
   //        1            if       x > 4
   //        0            if  -4 > x
@@ -518,12 +519,12 @@ Value sigmiod_seg3(HalContext* ctx, const Value& x) {
 
 }  // namespace
 
-Value f_sigmoid(HalContext* ctx, const Value& x) {
+Value f_sigmoid(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
 
-  switch (ctx->rt_config().sigmoid_mode()) {
+  switch (ctx->config().sigmoid_mode()) {
     case RuntimeConfig::SIGMOID_DEFAULT:
     case RuntimeConfig::SIGMOID_MM1: {
       return sigmiod_mm1(ctx, x);
diff --git a/libspu/kernel/hal/fxp_approx.h b/libspu/kernel/hal/fxp_approx.h
index 9b72ed34..84febd9a 100644
--- a/libspu/kernel/hal/fxp_approx.h
+++ b/libspu/kernel/hal/fxp_approx.h
@@ -14,43 +14,43 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 // !!please read [README.md] for api naming conventions.
 namespace spu::kernel::hal {
 namespace detail {
 
-Value log2_pade_approx(HalContext* ctx, const Value& x);
+Value log2_pade_approx(SPUContext* ctx, const Value& x);
 
-Value log_householder_approx(HalContext* ctx, const Value& x);
+Value log_householder_approx(SPUContext* ctx, const Value& x);
 
 // Works for range [-500, 2.1]
-Value exp_taylor_series(HalContext* ctx, const Value& x);
+Value exp_taylor_series(SPUContext* ctx, const Value& x);
 
-Value exp2_pade_approx(HalContext* ctx, const Value& x);
+Value exp2_pade_approx(SPUContext* ctx, const Value& x);
 
 // Works for range [-12.0, 18.0]
-Value exp_pade_approx(HalContext* ctx, const Value& x);
+Value exp_pade_approx(SPUContext* ctx, const Value& x);
 
 }  // namespace detail
 
-Value f_exp(HalContext* ctx, const Value& x);
+Value f_exp(SPUContext* ctx, const Value& x);
 
-Value f_log1p(HalContext* ctx, const Value& x);
+Value f_log1p(SPUContext* ctx, const Value& x);
 
-Value f_log(HalContext* ctx, const Value& x);
+Value f_log(SPUContext* ctx, const Value& x);
 
-Value f_log2(HalContext* ctx, const Value& x);
+Value f_log2(SPUContext* ctx, const Value& x);
 
-Value f_exp2(HalContext* ctx, const Value& x);
+Value f_exp2(SPUContext* ctx, const Value& x);
 
-Value f_tanh(HalContext* ctx, const Value& x);
+Value f_tanh(SPUContext* ctx, const Value& x);
 
-Value f_rsqrt(HalContext* ctx, const Value& x);
+Value f_rsqrt(SPUContext* ctx, const Value& x);
 
-Value f_sqrt(HalContext* ctx, const Value& x);
+Value f_sqrt(SPUContext* ctx, const Value& x);
 
-Value f_sigmoid(HalContext* ctx, const Value& x);
+Value f_sigmoid(SPUContext* ctx, const Value& x);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/fxp_base.cc b/libspu/kernel/hal/fxp_base.cc
index 9ab31b6f..1708992d 100644
--- a/libspu/kernel/hal/fxp_base.cc
+++ b/libspu/kernel/hal/fxp_base.cc
@@ -29,7 +29,7 @@ namespace detail {
 //
 // Coefficients should be ordered from the order 1 (linear) term first, ending
 // with the highest order term. (Constant is not included).
-Value f_polynomial(HalContext* ctx, const Value& x,
+Value f_polynomial(SPUContext* ctx, const Value& x,
                    const std::vector<Value>& coeffs) {
   SPU_TRACE_HAL_DISP(ctx, x);
   SPU_ENFORCE(x.isFxp());
@@ -52,7 +52,7 @@ Value f_polynomial(HalContext* ctx, const Value& x,
   return _trunc(ctx, res).asFxp();
 }
 
-Value highestOneBit(HalContext* ctx, const Value& x) {
+Value highestOneBit(SPUContext* ctx, const Value& x) {
   auto y = _prefix_or(ctx, x);
   auto y1 = _rshift(ctx, y, 1);
   return _xor(ctx, y, y1);
@@ -93,7 +93,7 @@ void hintNumberOfBits(const Value& a, size_t nbits) {
 //   return r * a * 2^{-m}
 //
 // Precision is decided by magic number, i.e 2.9142 and f.
-Value div_goldschmidt(HalContext* ctx, const Value& a, const Value& b) {
+Value div_goldschmidt(SPUContext* ctx, const Value& a, const Value& b) {
   SPU_TRACE_HAL_DISP(ctx, a, b);
 
   // We prefer  b_abs = b < 0 ? -b : b over b_abs = sign(b) * b
@@ -125,7 +125,7 @@ Value div_goldschmidt(HalContext* ctx, const Value& a, const Value& b) {
   auto r = w;
   auto e = f_sub(ctx, k1_, f_mul_with_sign(ctx, c, w, SignType::POSITIVE));
 
-  const size_t num_iters = ctx->rt_config().fxp_div_goldschmidt_iters();
+  const size_t num_iters = ctx->config().fxp_div_goldschmidt_iters();
   SPU_ENFORCE(num_iters != 0, "fxp_div_goldschmidt_iters should not be {}",
               num_iters);
 
@@ -144,7 +144,7 @@ Value div_goldschmidt(HalContext* ctx, const Value& a, const Value& b) {
   return _mux(ctx, is_negative, _negate(ctx, r), r).asFxp();
 }
 
-Value reciprocal_goldschmidt_positive(HalContext* ctx, const Value& b_abs) {
+Value reciprocal_goldschmidt_positive(SPUContext* ctx, const Value& b_abs) {
   auto b_msb = detail::highestOneBit(ctx, b_abs);
 
   // factor = 2^{2f-m} = 2^{f-m} * 2^f, the fixed point repr of 2^{f-m}
@@ -166,7 +166,7 @@ Value reciprocal_goldschmidt_positive(HalContext* ctx, const Value& b_abs) {
   auto r = w;
   auto e = f_sub(ctx, k1_, f_mul_with_sign(ctx, b_abs, w, SignType::POSITIVE));
 
-  const size_t num_iters = ctx->rt_config().fxp_div_goldschmidt_iters();
+  const size_t num_iters = ctx->config().fxp_div_goldschmidt_iters();
   SPU_ENFORCE(num_iters != 0, "fxp_div_goldschmidt_iters should not be {}",
               num_iters);
 
@@ -183,7 +183,7 @@ Value reciprocal_goldschmidt_positive(HalContext* ctx, const Value& b_abs) {
 
 // NOTE(junfeng): we have a seperate reciprocal_goldschmidt is to avoid
 // unnecessary f_mul for y initiation in div_goldschmidt.
-Value reciprocal_goldschmidt(HalContext* ctx, const Value& b) {
+Value reciprocal_goldschmidt(SPUContext* ctx, const Value& b) {
   SPU_TRACE_HAL_DISP(ctx, b);
 
   auto is_negative = _msb(ctx, b);
@@ -196,14 +196,14 @@ Value reciprocal_goldschmidt(HalContext* ctx, const Value& b) {
 
 }  // namespace detail
 
-Value f_negate(HalContext* ctx, const Value& x) {
+Value f_negate(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
   return _negate(ctx, x).asFxp();
 }
 
-Value f_abs(HalContext* ctx, const Value& x) {
+Value f_abs(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
@@ -212,7 +212,7 @@ Value f_abs(HalContext* ctx, const Value& x) {
   return _mul(ctx, sign, x).asFxp();
 }
 
-Value f_reciprocal(HalContext* ctx, const Value& x) {
+Value f_reciprocal(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
@@ -223,7 +223,7 @@ Value f_reciprocal(HalContext* ctx, const Value& x) {
   return detail::reciprocal_goldschmidt(ctx, x);
 }
 
-Value f_add(HalContext* ctx, const Value& x, const Value& y) {
+Value f_add(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   SPU_ENFORCE(x.isFxp());
@@ -232,7 +232,7 @@ Value f_add(HalContext* ctx, const Value& x, const Value& y) {
   return _add(ctx, x, y).asFxp();
 }
 
-Value f_sub(HalContext* ctx, const Value& x, const Value& y) {
+Value f_sub(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   SPU_ENFORCE(x.isFxp());
@@ -240,7 +240,7 @@ Value f_sub(HalContext* ctx, const Value& x, const Value& y) {
   return f_add(ctx, x, f_negate(ctx, y));
 }
 
-Value f_mul(HalContext* ctx, const Value& x, const Value& y) {
+Value f_mul(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   SPU_ENFORCE(x.isFxp());
@@ -248,7 +248,7 @@ Value f_mul(HalContext* ctx, const Value& x, const Value& y) {
   return _trunc(ctx, _mul(ctx, x, y)).asFxp();
 }
 
-Value f_mul_with_sign(HalContext* ctx, const Value& x, const Value& y,
+Value f_mul_with_sign(SPUContext* ctx, const Value& x, const Value& y,
                       SignType sign) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
@@ -273,7 +273,7 @@ Value f_mul_with_sign(HalContext* ctx, const Value& x, const Value& y,
   }
 }
 
-Value f_mmul(HalContext* ctx, const Value& x, const Value& y) {
+Value f_mmul(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   SPU_ENFORCE(x.isFxp());
@@ -282,7 +282,7 @@ Value f_mmul(HalContext* ctx, const Value& x, const Value& y) {
   return _trunc(ctx, _mmul(ctx, x, y)).asFxp();
 }
 
-Value f_conv2d(HalContext* ctx, const Value& x, const Value& y,
+Value f_conv2d(SPUContext* ctx, const Value& x, const Value& y,
                absl::Span<const int64_t> window_strides,
                absl::Span<const int64_t> result_shape) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
@@ -293,7 +293,7 @@ Value f_conv2d(HalContext* ctx, const Value& x, const Value& y,
   return _trunc(ctx, _conv2d(ctx, x, y, window_strides, result_shape)).asFxp();
 }
 
-Value f_div(HalContext* ctx, const Value& x, const Value& y) {
+Value f_div(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   SPU_ENFORCE(x.isFxp());
@@ -306,7 +306,7 @@ Value f_div(HalContext* ctx, const Value& x, const Value& y) {
   return detail::div_goldschmidt(ctx, x, y);
 }
 
-Value f_equal(HalContext* ctx, const Value& x, const Value& y) {
+Value f_equal(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   SPU_ENFORCE(x.isFxp());
@@ -315,7 +315,7 @@ Value f_equal(HalContext* ctx, const Value& x, const Value& y) {
   return _equal(ctx, x, y).setDtype(DT_I1);
 }
 
-Value f_less(HalContext* ctx, const Value& x, const Value& y) {
+Value f_less(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   SPU_ENFORCE(x.isFxp());
@@ -324,7 +324,7 @@ Value f_less(HalContext* ctx, const Value& x, const Value& y) {
   return _less(ctx, x, y).setDtype(DT_I1);
 }
 
-Value f_square(HalContext* ctx, const Value& x) {
+Value f_square(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
@@ -335,7 +335,7 @@ Value f_square(HalContext* ctx, const Value& x) {
       .asFxp();
 }
 
-Value f_floor(HalContext* ctx, const Value& x) {
+Value f_floor(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
@@ -344,7 +344,7 @@ Value f_floor(HalContext* ctx, const Value& x) {
   return _lshift(ctx, _arshift(ctx, x, fbits), fbits).asFxp();
 }
 
-Value f_ceil(HalContext* ctx, const Value& x) {
+Value f_ceil(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
diff --git a/libspu/kernel/hal/fxp_base.h b/libspu/kernel/hal/fxp_base.h
index 8e7563a8..44029b22 100644
--- a/libspu/kernel/hal/fxp_base.h
+++ b/libspu/kernel/hal/fxp_base.h
@@ -14,9 +14,9 @@
 
 #pragma once
 
+#include "libspu/core/context.h"
 #include "libspu/core/pt_buffer_view.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/value.h"
 
 // !!please read [README.md] for api naming conventions.
 namespace spu::kernel::hal {
@@ -24,17 +24,17 @@ namespace detail {
 
 // Extract the most significant bit. see
 // https://docs.oracle.com/javase/7/docs/api/java/lang/Integer.html#highestOneBit(int)
-Value highestOneBit(HalContext* ctx, const Value& x);
+Value highestOneBit(SPUContext* ctx, const Value& x);
 
 void hintNumberOfBits(const Value& a, size_t nbits);
 
-Value div_goldschmidt(HalContext* ctx, const Value& a, const Value& b);
+Value div_goldschmidt(SPUContext* ctx, const Value& a, const Value& b);
 
-Value reciprocal_goldschmidt_positive(HalContext* ctx, const Value& b_abs);
+Value reciprocal_goldschmidt_positive(SPUContext* ctx, const Value& b_abs);
 
-Value reciprocal_goldschmidt(HalContext* ctx, const Value& b);
+Value reciprocal_goldschmidt(SPUContext* ctx, const Value& b);
 
-Value f_polynomial(HalContext* ctx, const Value& x,
+Value f_polynomial(SPUContext* ctx, const Value& x,
                    const std::vector<Value>& coeffs);
 
 }  // namespace detail
@@ -46,37 +46,37 @@ enum class SignType {
   NEGATIVE,
 };
 
-Value f_negate(HalContext* ctx, const Value& x);
+Value f_negate(SPUContext* ctx, const Value& x);
 
-Value f_abs(HalContext* ctx, const Value& x);
+Value f_abs(SPUContext* ctx, const Value& x);
 
-Value f_reciprocal(HalContext* ctx, const Value& x);
+Value f_reciprocal(SPUContext* ctx, const Value& x);
 
-Value f_add(HalContext* ctx, const Value& x, const Value& y);
+Value f_add(SPUContext* ctx, const Value& x, const Value& y);
 
-Value f_sub(HalContext* ctx, const Value& x, const Value& y);
+Value f_sub(SPUContext* ctx, const Value& x, const Value& y);
 
-Value f_mul(HalContext* ctx, const Value& x, const Value& y);
+Value f_mul(SPUContext* ctx, const Value& x, const Value& y);
 
-Value f_mul_with_sign(HalContext* ctx, const Value& x, const Value& y,
+Value f_mul_with_sign(SPUContext* ctx, const Value& x, const Value& y,
                       SignType sign = SignType::UNKNOWN);
 
-Value f_mmul(HalContext* ctx, const Value& x, const Value& y);
+Value f_mmul(SPUContext* ctx, const Value& x, const Value& y);
 
-Value f_conv2d(HalContext* ctx, const Value& x, const Value& y,
+Value f_conv2d(SPUContext* ctx, const Value& x, const Value& y,
                absl::Span<const int64_t> window_strides,
                absl::Span<const int64_t> result_shape);
 
-Value f_div(HalContext* ctx, const Value& x, const Value& y);
+Value f_div(SPUContext* ctx, const Value& x, const Value& y);
 
-Value f_equal(HalContext* ctx, const Value& x, const Value& y);
+Value f_equal(SPUContext* ctx, const Value& x, const Value& y);
 
-Value f_less(HalContext* ctx, const Value& x, const Value& y);
+Value f_less(SPUContext* ctx, const Value& x, const Value& y);
 
-Value f_square(HalContext* ctx, const Value& x);
+Value f_square(SPUContext* ctx, const Value& x);
 
-Value f_floor(HalContext* ctx, const Value& x);
+Value f_floor(SPUContext* ctx, const Value& x);
 
-Value f_ceil(HalContext* ctx, const Value& x);
+Value f_ceil(SPUContext* ctx, const Value& x);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/fxp_cleartext.cc b/libspu/kernel/hal/fxp_cleartext.cc
index 4382fe12..e3c2fa0e 100644
--- a/libspu/kernel/hal/fxp_cleartext.cc
+++ b/libspu/kernel/hal/fxp_cleartext.cc
@@ -22,7 +22,7 @@ namespace spu::kernel::hal {
 namespace {
 
 template <typename FN>
-Value applyFloatingPointFn(HalContext* ctx, const Value& in, FN&& fn) {
+Value applyFloatingPointFn(SPUContext* ctx, const Value& in, FN&& fn) {
   SPU_TRACE_HAL_DISP(ctx, in);
   SPU_ENFORCE(in.isPublic(), "expected public, got {}", in.storage_type());
   SPU_ENFORCE(in.dtype() == DT_FXP, "expected fxp, got={}", in.dtype());
@@ -47,7 +47,7 @@ Value applyFloatingPointFn(HalContext* ctx, const Value& in, FN&& fn) {
 }
 
 template <typename FN>
-Value applyFloatingPointFn(HalContext* ctx, const Value& x, const Value& y,
+Value applyFloatingPointFn(SPUContext* ctx, const Value& x, const Value& y,
                            FN&& fn) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.isPublic() && y.isPublic(), "expect public, got {}, {}",
@@ -79,23 +79,23 @@ Value applyFloatingPointFn(HalContext* ctx, const Value& x, const Value& y,
 
 }  // namespace
 
-Value f_reciprocal_p(HalContext* ctx, const Value& in) {
+Value f_reciprocal_p(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   return applyFloatingPointFn(ctx, in, [](float x) { return 1.0 / x; });
 }
 
-Value f_log_p(HalContext* ctx, const Value& in) {
+Value f_log_p(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
   return applyFloatingPointFn(ctx, in, [](float x) { return std::log(x); });
 }
 
-Value f_exp_p(HalContext* ctx, const Value& in) {
+Value f_exp_p(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
   return applyFloatingPointFn(ctx, in, [](float x) { return std::exp(x); });
 }
 
-Value f_div_p(HalContext* ctx, const Value& x, const Value& y) {
+Value f_div_p(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   return applyFloatingPointFn(ctx, x, y,
                               [](float a, float b) { return a / b; });
diff --git a/libspu/kernel/hal/fxp_cleartext.h b/libspu/kernel/hal/fxp_cleartext.h
index 2399c3c3..3a09aa47 100644
--- a/libspu/kernel/hal/fxp_cleartext.h
+++ b/libspu/kernel/hal/fxp_cleartext.h
@@ -14,8 +14,8 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 // !!please read [README.md] for api naming conventions.
 namespace spu::kernel::hal {
@@ -23,12 +23,12 @@ namespace spu::kernel::hal {
 // This module decode public fxp back to flp and use std math function for
 // evaluation.
 
-Value f_reciprocal_p(HalContext* ctx, const Value& in);
+Value f_reciprocal_p(SPUContext* ctx, const Value& in);
 
-Value f_log_p(HalContext* ctx, const Value& in);
+Value f_log_p(SPUContext* ctx, const Value& in);
 
-Value f_exp_p(HalContext* ctx, const Value& in);
+Value f_exp_p(SPUContext* ctx, const Value& in);
 
-Value f_div_p(HalContext* ctx, const Value& x, const Value& y);
+Value f_div_p(SPUContext* ctx, const Value& x, const Value& y);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/integer.cc b/libspu/kernel/hal/integer.cc
index 69edfbdb..1b094361 100644
--- a/libspu/kernel/hal/integer.cc
+++ b/libspu/kernel/hal/integer.cc
@@ -25,7 +25,7 @@ namespace spu::kernel::hal {
   SPU_ENFORCE((Y).isInt(), "expect rhs int, got {]", (X).dtype());
 
 #define DEF_UNARY_OP(Name, Fn2K)                             \
-  Value Name(HalContext* ctx, const Value& x) {              \
+  Value Name(SPUContext* ctx, const Value& x) {              \
     SPU_TRACE_HAL_LEAF(ctx, x);                              \
                                                              \
     SPU_ENFORCE(x.isInt(), "expect Int, got {]", x.dtype()); \
@@ -38,7 +38,7 @@ DEF_UNARY_OP(i_negate, _negate)
 #undef DEF_UNARY_OP
 
 #define DEF_BINARY_OP(Name, Fn2K)                               \
-  Value Name(HalContext* ctx, const Value& x, const Value& y) { \
+  Value Name(SPUContext* ctx, const Value& x, const Value& y) { \
     SPU_TRACE_HAL_LEAF(ctx, x, y);                              \
     ENSURE_INT_AND_DTYPE_MATCH(x, y);                           \
     return Fn2K(ctx, x, y).setDtype(x.dtype());                 \
@@ -48,7 +48,7 @@ DEF_BINARY_OP(i_add, _add)
 DEF_BINARY_OP(i_mul, _mul)
 DEF_BINARY_OP(i_mmul, _mmul)
 
-Value i_less(HalContext* ctx, const Value& x, const Value& y) {
+Value i_less(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
   ENSURE_INT_AND_DTYPE_MATCH(x, y);
 
@@ -57,7 +57,7 @@ Value i_less(HalContext* ctx, const Value& x, const Value& y) {
 
 #undef DEF_BINARY_OP
 
-Value i_abs(HalContext* ctx, const Value& x) {
+Value i_abs(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isInt());
@@ -66,7 +66,7 @@ Value i_abs(HalContext* ctx, const Value& x) {
   return _mul(ctx, _sign(ctx, x), x).setDtype(x.dtype());
 }
 
-Value i_equal(HalContext* ctx, const Value& x, const Value& y) {
+Value i_equal(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   SPU_ENFORCE(x.isInt());
@@ -75,7 +75,7 @@ Value i_equal(HalContext* ctx, const Value& x, const Value& y) {
   return _equal(ctx, x, y).setDtype(DT_I1);
 }
 
-Value i_sub(HalContext* ctx, const Value& x, const Value& y) {
+Value i_sub(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   SPU_ENFORCE(x.isInt());
@@ -83,7 +83,7 @@ Value i_sub(HalContext* ctx, const Value& x, const Value& y) {
   return i_add(ctx, x, i_negate(ctx, y));
 }
 
-Value i_conv2d(HalContext* ctx, const Value& x, const Value& y,
+Value i_conv2d(SPUContext* ctx, const Value& x, const Value& y,
                absl::Span<const int64_t> window_strides,
                absl::Span<const int64_t> result_shape) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
diff --git a/libspu/kernel/hal/integer.h b/libspu/kernel/hal/integer.h
index 93a27979..64f8babc 100644
--- a/libspu/kernel/hal/integer.h
+++ b/libspu/kernel/hal/integer.h
@@ -14,8 +14,8 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hal {
 
@@ -25,24 +25,24 @@ namespace spu::kernel::hal {
 // security semantics, it dispatches functionality by value's security type to
 // the underline mpc module.
 
-Value i_negate(HalContext* ctx, const Value& x);
+Value i_negate(SPUContext* ctx, const Value& x);
 
-Value i_abs(HalContext* ctx, const Value& x);
+Value i_abs(SPUContext* ctx, const Value& x);
 
-Value i_add(HalContext* ctx, const Value& x, const Value& y);
+Value i_add(SPUContext* ctx, const Value& x, const Value& y);
 
-Value i_sub(HalContext* ctx, const Value& x, const Value& y);
+Value i_sub(SPUContext* ctx, const Value& x, const Value& y);
 
-Value i_mul(HalContext* ctx, const Value& x, const Value& y);
+Value i_mul(SPUContext* ctx, const Value& x, const Value& y);
 
-Value i_mmul(HalContext* ctx, const Value& x, const Value& y);
+Value i_mmul(SPUContext* ctx, const Value& x, const Value& y);
 
-Value i_conv2d(HalContext* ctx, const Value& x, const Value& y,
+Value i_conv2d(SPUContext* ctx, const Value& x, const Value& y,
                absl::Span<const int64_t> window_strides,
                absl::Span<const int64_t> result_shape);
 
-Value i_equal(HalContext* ctx, const Value& x, const Value& y);
+Value i_equal(SPUContext* ctx, const Value& x, const Value& y);
 
-Value i_less(HalContext* ctx, const Value& x, const Value& y);
+Value i_less(SPUContext* ctx, const Value& x, const Value& y);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/polymorphic.cc b/libspu/kernel/hal/polymorphic.cc
index b4021d3f..a2990d9f 100644
--- a/libspu/kernel/hal/polymorphic.cc
+++ b/libspu/kernel/hal/polymorphic.cc
@@ -17,10 +17,10 @@
 #include "fmt/format.h"
 #include "fmt/ostream.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/encoding.h"  // for bitcast
 #include "libspu/core/prelude.h"
 #include "libspu/core/trace.h"
-#include "libspu/kernel/context.h"
 #include "libspu/kernel/hal/constants.h"
 #include "libspu/kernel/hal/fxp.h"
 #include "libspu/kernel/hal/integer.h"
@@ -32,8 +32,8 @@
 namespace spu::kernel::hal {
 namespace {
 
-using UnaryOp = Value(HalContext*, const Value&);
-using BinaryOp = Value(HalContext*, const Value&, const Value&);
+using UnaryOp = Value(SPUContext*, const Value&);
+using BinaryOp = Value(SPUContext*, const Value&, const Value&);
 
 DataType common_dtype(DataType lhs, DataType rhs) {
   if (lhs == rhs) {
@@ -43,7 +43,7 @@ DataType common_dtype(DataType lhs, DataType rhs) {
 }
 
 template <BinaryOp* FnFxp, BinaryOp* FnInt>
-Value dtypeBinaryDispatch(std::string_view op_name, HalContext* ctx,
+Value dtypeBinaryDispatch(std::string_view op_name, SPUContext* ctx,
                           const Value& x, const Value& y) {
   // Promote int to fxp if mismatch.
   if (x.isInt() && y.isInt()) {
@@ -62,7 +62,7 @@ Value dtypeBinaryDispatch(std::string_view op_name, HalContext* ctx,
 }
 
 template <UnaryOp* FnFxp, UnaryOp* FnInt>
-Value dtypeUnaryDispatch(std::string_view op_name, HalContext* ctx,
+Value dtypeUnaryDispatch(std::string_view op_name, SPUContext* ctx,
                          const Value& x) {
   // Promote int to fxp if mismatch.
   if (x.isInt()) {
@@ -80,27 +80,27 @@ bool isCrossIntFxp(const Value& x, const Value& y) {
 
 }  // namespace
 
-Value add(HalContext* ctx, const Value& x, const Value& y) {
+Value add(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   return dtypeBinaryDispatch<f_add, i_add>("add", ctx, x, y);
 }
 
-Value sub(HalContext* ctx, const Value& x, const Value& y) {
+Value sub(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   return dtypeBinaryDispatch<f_sub, i_sub>("sub", ctx, x, y);
 }
 
-Value mixed_mul(HalContext* ctx, const Value& x, const Value& y) {
+Value mixed_mul(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
   return _mul(ctx, x, y).asFxp();
 }
 
-Value mixed_mmul(HalContext* ctx, const Value& x, const Value& y) {
+Value mixed_mmul(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
   return _mmul(ctx, x, y).asFxp();
 }
 
-Value mul(HalContext* ctx, const Value& x, const Value& y) {
+Value mul(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   // fast dispatch, avoid truncation cost
   if (isCrossIntFxp(x, y)) {
@@ -110,7 +110,7 @@ Value mul(HalContext* ctx, const Value& x, const Value& y) {
   return dtypeBinaryDispatch<f_mul, i_mul>("mul", ctx, x, y);
 }
 
-Value matmul(HalContext* ctx, const Value& x, const Value& y) {
+Value matmul(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   // fast dispatch, avoid truncation cost
   if (isCrossIntFxp(x, y)) {
@@ -120,7 +120,7 @@ Value matmul(HalContext* ctx, const Value& x, const Value& y) {
   return dtypeBinaryDispatch<f_mmul, i_mmul>("mmul", ctx, x, y);
 }
 
-Value logical_not(HalContext* ctx, const Value& in) {
+Value logical_not(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_LEAF(ctx, in);
 
   auto _k1 = _constant(ctx, 1, in.shape());
@@ -133,28 +133,28 @@ Value logical_not(HalContext* ctx, const Value& in) {
   }
 }
 
-Value equal(HalContext* ctx, const Value& x, const Value& y) {
+Value equal(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.shape() == y.shape(), "x = {}, y = {}", x, y);
 
   return dtypeBinaryDispatch<f_equal, i_equal>("equal", ctx, x, y);
 }
 
-Value not_equal(HalContext* ctx, const Value& x, const Value& y) {
+Value not_equal(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.shape() == y.shape());
 
   return logical_not(ctx, equal(ctx, x, y));
 }
 
-Value less(HalContext* ctx, const Value& x, const Value& y) {
+Value less(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.shape() == y.shape());
 
   return dtypeBinaryDispatch<f_less, i_less>("less", ctx, x, y);
 }
 
-Value less_equal(HalContext* ctx, const Value& x, const Value& y) {
+Value less_equal(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.shape() == y.shape());
 
@@ -162,14 +162,14 @@ Value less_equal(HalContext* ctx, const Value& x, const Value& y) {
   return logical_not(ctx, greater(ctx, x, y));
 }
 
-Value greater(HalContext* ctx, const Value& x, const Value& y) {
+Value greater(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.shape() == y.shape());
 
   return less(ctx, y, x);
 }
 
-Value greater_equal(HalContext* ctx, const Value& x, const Value& y) {
+Value greater_equal(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.shape() == y.shape());
 
@@ -177,25 +177,25 @@ Value greater_equal(HalContext* ctx, const Value& x, const Value& y) {
   return logical_not(ctx, less(ctx, x, y));
 }
 
-Value negate(HalContext* ctx, const Value& x) {
+Value negate(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
   return dtypeUnaryDispatch<f_negate, i_negate>("negate", ctx, x);
 }
 
-Value abs(HalContext* ctx, const Value& x) {
+Value abs(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
   return dtypeUnaryDispatch<f_abs, i_abs>("abs", ctx, x);
 }
 
-Value exp(HalContext* ctx, const Value& in) {
+Value exp(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   return f_exp(ctx, dtype_cast(ctx, in, DT_FXP));
 }
 
-Value select(HalContext* ctx, const Value& pred, const Value& a,
+Value select(SPUContext* ctx, const Value& pred, const Value& a,
              const Value& b) {
   SPU_TRACE_HAL_DISP(ctx, pred, a, b);
 
@@ -209,7 +209,7 @@ Value select(HalContext* ctx, const Value& pred, const Value& a,
   return _mux(ctx, pred, a, b).setDtype(a.dtype());
 }
 
-Value bitwise_and(HalContext* ctx, const Value& x, const Value& y) {
+Value bitwise_and(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
 
   SPU_ENFORCE(x.isInt() && y.isInt());
@@ -218,7 +218,7 @@ Value bitwise_and(HalContext* ctx, const Value& x, const Value& y) {
   return _and(ctx, x, y).setDtype(x.dtype());
 }
 
-Value bitwise_xor(HalContext* ctx, const Value& x, const Value& y) {
+Value bitwise_xor(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.isInt() && y.isInt());
   SPU_ENFORCE(x.shape() == y.shape());
@@ -226,7 +226,7 @@ Value bitwise_xor(HalContext* ctx, const Value& x, const Value& y) {
   return _xor(ctx, x, y).setDtype(x.dtype());
 }
 
-Value bitwise_or(HalContext* ctx, const Value& x, const Value& y) {
+Value bitwise_or(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.isInt() && y.isInt());
   SPU_ENFORCE(x.shape() == y.shape());
@@ -234,39 +234,39 @@ Value bitwise_or(HalContext* ctx, const Value& x, const Value& y) {
   return _or(ctx, x, y).setDtype(x.dtype());
 }
 
-Value bitwise_not(HalContext* ctx, const Value& in) {
+Value bitwise_not(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   return _not(ctx, in).setDtype(in.dtype());
 }
 
-Value logistic(HalContext* ctx, const Value& in) {
+Value logistic(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   SPU_ENFORCE(in.isFxp());
   return f_sigmoid(ctx, in);
 }
 
-Value log(HalContext* ctx, const Value& in) {
+Value log(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   return f_log(ctx, dtype_cast(ctx, in, DT_FXP));
 }
 
-Value log1p(HalContext* ctx, const Value& in) {
+Value log1p(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   return f_log1p(ctx, dtype_cast(ctx, in, DT_FXP));
 }
 
-Value reciprocal(HalContext* ctx, const Value& in) {
+Value reciprocal(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
   SPU_ENFORCE(in.isFxp());
 
   return f_reciprocal(ctx, in);
 }
 
-Value floor(HalContext* ctx, const Value& in) {
+Value floor(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   SPU_ENFORCE(in.isFxp());
@@ -274,7 +274,7 @@ Value floor(HalContext* ctx, const Value& in) {
   return f_floor(ctx, in);
 }
 
-Value ceil(HalContext* ctx, const Value& in) {
+Value ceil(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   SPU_ENFORCE(in.isFxp());
@@ -282,7 +282,7 @@ Value ceil(HalContext* ctx, const Value& in) {
   return f_ceil(ctx, in);
 }
 
-Value max(HalContext* ctx, const Value& x, const Value& y) {
+Value max(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
 
   SPU_ENFORCE(x.dtype() == y.dtype());
@@ -290,7 +290,7 @@ Value max(HalContext* ctx, const Value& x, const Value& y) {
   return select(ctx, greater(ctx, x, y), x, y);
 }
 
-Value min(HalContext* ctx, const Value& x, const Value& y) {
+Value min(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
 
   SPU_ENFORCE(x.dtype() == y.dtype());
@@ -298,13 +298,13 @@ Value min(HalContext* ctx, const Value& x, const Value& y) {
   return select(ctx, less(ctx, x, y), x, y);
 }
 
-Value power(HalContext* ctx, const Value& x, const Value& y) {
+Value power(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
   // x^y = e^(y*ln(x))
   return exp(ctx, mul(ctx, y, log(ctx, x)));
 }
 
-Value idiv(HalContext* ctx, const Value& x, const Value& y) {
+Value idiv(SPUContext* ctx, const Value& x, const Value& y) {
   auto sign_x = sign(ctx, x);
   auto sign_y = sign(ctx, y);
 
@@ -335,7 +335,7 @@ Value idiv(HalContext* ctx, const Value& x, const Value& y) {
   return mul(ctx, q, mul(ctx, sign_x, sign_y));
 }
 
-Value div(HalContext* ctx, const Value& x, const Value& y) {
+Value div(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
 
   if (x.isInt() && y.isInt()) {
@@ -355,7 +355,7 @@ Value div(HalContext* ctx, const Value& x, const Value& y) {
   return res_f;
 }
 
-Value clamp(HalContext* ctx, const Value& x, const Value& minv,
+Value clamp(SPUContext* ctx, const Value& x, const Value& minv,
             const Value& maxv) {
   SPU_TRACE_HAL_DISP(ctx, x, minv, maxv);
 
@@ -366,44 +366,44 @@ Value clamp(HalContext* ctx, const Value& x, const Value& minv,
   return min(ctx, max(ctx, minv, x), maxv);
 }
 
-Value bitcast(HalContext* ctx, const Value& x, DataType dtype) {
+Value bitcast(SPUContext* ctx, const Value& x, DataType dtype) {
   SPU_TRACE_HAL_DISP(ctx, x, dtype);
 
   // FIXME(jint) should we directly use fixed point binary expr for bitcast?
   return Value(x.data().clone(), dtype);
 }
 
-Value left_shift(HalContext* ctx, const Value& x, size_t bits) {
+Value left_shift(SPUContext* ctx, const Value& x, size_t bits) {
   SPU_TRACE_HAL_DISP(ctx, x, bits);
 
   return _lshift(ctx, x, bits).setDtype(x.dtype());
 }
 
-Value right_shift_logical(HalContext* ctx, const Value& x, size_t bits) {
+Value right_shift_logical(SPUContext* ctx, const Value& x, size_t bits) {
   SPU_TRACE_HAL_DISP(ctx, x, bits);
 
   return _rshift(ctx, x, bits).setDtype(x.dtype());
 }
 
-Value right_shift_arithmetic(HalContext* ctx, const Value& x, size_t bits) {
+Value right_shift_arithmetic(SPUContext* ctx, const Value& x, size_t bits) {
   SPU_TRACE_HAL_DISP(ctx, x, bits);
 
   return _arshift(ctx, x, bits).setDtype(x.dtype());
 }
 
-Value log2(HalContext* ctx, const Value& in) {
+Value log2(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   return f_log2(ctx, dtype_cast(ctx, in, DT_FXP));
 }
 
-Value exp2(HalContext* ctx, const Value& x) {
+Value exp2(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
   return f_exp2(ctx, dtype_cast(ctx, x, DT_FXP));
 }
 
-Value tanh(HalContext* ctx, const Value& x) {
+Value tanh(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
   // For tanh inputs beyond [-3, 3], result is infinitely close to -1, 1
@@ -416,25 +416,25 @@ Value tanh(HalContext* ctx, const Value& x) {
   return f_tanh(ctx, normalized_x);
 }
 
-Value rsqrt(HalContext* ctx, const Value& x) {
+Value rsqrt(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
   return f_rsqrt(ctx, dtype_cast(ctx, x, DT_FXP));
 }
 
-Value sqrt(HalContext* ctx, const Value& x) {
+Value sqrt(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
   return f_sqrt(ctx, dtype_cast(ctx, x, DT_FXP));
 }
 
-Value sign(HalContext* ctx, const Value& x) {
+Value sign(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
   return _sign(ctx, x).setDtype(DT_I8);
 }
 
-Value conv2d(HalContext* ctx, const Value& x, const Value& y,
+Value conv2d(SPUContext* ctx, const Value& x, const Value& y,
              absl::Span<const int64_t> window_strides,
              absl::Span<const int64_t> result_shape) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
diff --git a/libspu/kernel/hal/polymorphic.h b/libspu/kernel/hal/polymorphic.h
index 9b7ab33c..a3b98787 100644
--- a/libspu/kernel/hal/polymorphic.h
+++ b/libspu/kernel/hal/polymorphic.h
@@ -14,196 +14,196 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hal {
 
 /// the element-wise absolute value function
 // @param in, the value to negate
-Value abs(HalContext* ctx, const Value& x);
+Value abs(SPUContext* ctx, const Value& x);
 
 /// general element-wise add operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value add(HalContext* ctx, const Value& x, const Value& y);
+Value add(SPUContext* ctx, const Value& x, const Value& y);
 
 /// general element-wise bitwise and operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value bitwise_and(HalContext* ctx, const Value& x, const Value& y);
+Value bitwise_and(SPUContext* ctx, const Value& x, const Value& y);
 
 /// general element-wise bitwise xor operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value bitwise_xor(HalContext* ctx, const Value& x, const Value& y);
+Value bitwise_xor(SPUContext* ctx, const Value& x, const Value& y);
 
 /// general element-wise bitwise or operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value bitwise_or(HalContext* ctx, const Value& x, const Value& y);
+Value bitwise_or(SPUContext* ctx, const Value& x, const Value& y);
 
 /// see numpy.bitwise_not(in)
 // @param in, the input parameter
-Value bitwise_not(HalContext* ctx, const Value& in);
+Value bitwise_not(SPUContext* ctx, const Value& in);
 
 /// matrix production operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value matmul(HalContext* ctx, const Value& x, const Value& y);
+Value matmul(SPUContext* ctx, const Value& x, const Value& y);
 
 /// 2-dimensional convolution operator
 // @param x, the input tensor
 // @param y, the kernel weight
 // @param window_strides, sized-2 window strides
 // @param result_shape, output shape
-Value conv2d(HalContext* ctx, const Value& x, const Value& y,
+Value conv2d(SPUContext* ctx, const Value& x, const Value& y,
              absl::Span<const int64_t> window_strides,
              absl::Span<const int64_t> result_shape);
 
 /// general element-wise bitwise equal operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value equal(HalContext* ctx, const Value& x, const Value& y);
+Value equal(SPUContext* ctx, const Value& x, const Value& y);
 
 /// element-wise natural exponential x -> e^x
 // @param in, the input value
-Value exp(HalContext* ctx, const Value& in);
+Value exp(SPUContext* ctx, const Value& in);
 
 /// element-wise floor
 // @param in, the input value
-Value floor(HalContext* ctx, const Value& in);
+Value floor(SPUContext* ctx, const Value& in);
 
 /// element-wise ceil
 // @param in, the input value
-Value ceil(HalContext* ctx, const Value& in);
+Value ceil(SPUContext* ctx, const Value& in);
 
 /// general element-wise bitwise greater operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value greater(HalContext* ctx, const Value& x, const Value& y);
+Value greater(SPUContext* ctx, const Value& x, const Value& y);
 
 /// general element-wise bitwise greater or equal operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value greater_equal(HalContext* ctx, const Value& x, const Value& y);
+Value greater_equal(SPUContext* ctx, const Value& x, const Value& y);
 
 /// general element-wise bitwise less operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value less(HalContext* ctx, const Value& x, const Value& y);
+Value less(SPUContext* ctx, const Value& x, const Value& y);
 
 /// general element-wise bitwise less or equal operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value less_equal(HalContext* ctx, const Value& x, const Value& y);
+Value less_equal(SPUContext* ctx, const Value& x, const Value& y);
 
 /// the element-wise natural logarithm
 // @param in, the param
-Value log(HalContext* ctx, const Value& in);
+Value log(SPUContext* ctx, const Value& in);
 
 /// the element-wise natural logarithm of (1 + x)
 // @param in, the param
-Value log1p(HalContext* ctx, const Value& in);
+Value log1p(SPUContext* ctx, const Value& in);
 
 /// see numpy.logical_not(in)
 // @param in, requires integer one or zero
-Value logical_not(HalContext* ctx, const Value& in);
+Value logical_not(SPUContext* ctx, const Value& in);
 
 /// the element-wise sigmoid function
 // @param in, the param
-Value logistic(HalContext* ctx, const Value& in);
+Value logistic(SPUContext* ctx, const Value& in);
 
 /// element-wise maximum
 // @param x, first input value
 // @param y, second input value
-Value max(HalContext* ctx, const Value& x, const Value& y);
+Value max(SPUContext* ctx, const Value& x, const Value& y);
 
 /// element-wise minimum
 // @param x, first input value
 // @param y, second input value
-Value min(HalContext* ctx, const Value& x, const Value& y);
+Value min(SPUContext* ctx, const Value& x, const Value& y);
 
 /// general element-wise multiply operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value mul(HalContext* ctx, const Value& x, const Value& y);
+Value mul(SPUContext* ctx, const Value& x, const Value& y);
 
-Value div(HalContext* ctx, const Value& x, const Value& y);
+Value div(SPUContext* ctx, const Value& x, const Value& y);
 
 /// see numpy.negate(in)
 // @param in, the value to negate
-Value negate(HalContext* ctx, const Value& x);
+Value negate(SPUContext* ctx, const Value& x);
 
 /// general element-wise bitwise equal operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value not_equal(HalContext* ctx, const Value& x, const Value& y);
+Value not_equal(SPUContext* ctx, const Value& x, const Value& y);
 
 /// element-wise power x ^ y
 // @param x, first input value, must be positive at this moment.
 // @param y, second input value
 // FIXME(junfeng): fix negative x.
-Value power(HalContext* ctx, const Value& x, const Value& y);
+Value power(SPUContext* ctx, const Value& x, const Value& y);
 
 /// the element-wise reciprocal function
 // @param in, the param
-Value reciprocal(HalContext* ctx, const Value& in);
+Value reciprocal(SPUContext* ctx, const Value& in);
 
 /// see numpy.select
 // @param pred, the predicate, requires integer zero or one
 // @param a, the first param
 // @param b, the second param
-Value select(HalContext* ctx, const Value& pred, const Value& a,
+Value select(SPUContext* ctx, const Value& pred, const Value& a,
              const Value& b);
 
 /// general element-wise subtract operator
 // @param x, the first parameter
 // @param y, the second parameter
-Value sub(HalContext* ctx, const Value& x, const Value& y);
+Value sub(SPUContext* ctx, const Value& x, const Value& y);
 
 /// general element-wise clamp operator
 // @param x, the first parameter
 // @param min, the second parameter
 // @param max, the third parameter
-Value clamp(HalContext* ctx, const Value& x, const Value& min,
+Value clamp(SPUContext* ctx, const Value& x, const Value& min,
             const Value& max);
 
 /// element-wise bitcast (reinterpret_cast)
 // @param x, first input value
 // @param dtype, second input value
-Value bitcast(HalContext* ctx, const Value& x, DataType dtype);
+Value bitcast(SPUContext* ctx, const Value& x, DataType dtype);
 
-Value left_shift(HalContext* ctx, const Value& x, size_t bits);
+Value left_shift(SPUContext* ctx, const Value& x, size_t bits);
 
-Value right_shift_logical(HalContext* ctx, const Value& x, size_t bits);
+Value right_shift_logical(SPUContext* ctx, const Value& x, size_t bits);
 
-Value right_shift_arithmetic(HalContext* ctx, const Value& x, size_t bits);
+Value right_shift_arithmetic(SPUContext* ctx, const Value& x, size_t bits);
 
-Value popcount(HalContext* ctx, const Value& x);
+Value popcount(SPUContext* ctx, const Value& x);
 
 /// the element-wise base-2 logarithm of x
 // @param in, should be positive, or the result is implementation defined.
-Value log2(HalContext* ctx, const Value& in);
+Value log2(SPUContext* ctx, const Value& in);
 
 /// element-wise 2 to the power x, i.e. x -> 2^x
 // @param in, the input value
-Value exp2(HalContext* ctx, const Value& x);
+Value exp2(SPUContext* ctx, const Value& x);
 
 /// element-wise hyperbolic tangent, i.e. x -> tanh(x)
 // @param in, the input value
-Value tanh(HalContext* ctx, const Value& x);
+Value tanh(SPUContext* ctx, const Value& x);
 
 /// element-wise reciprocal of square root operation, i.e. x - > 1.0 / sqrt(x)
 // @param in, the input value
-Value rsqrt(HalContext* ctx, const Value& x);
+Value rsqrt(SPUContext* ctx, const Value& x);
 
 /// element-wise square root operation.
 // @param in, the input value
-Value sqrt(HalContext* ctx, const Value& x);
+Value sqrt(SPUContext* ctx, const Value& x);
 
 /// element-wise sign operation
 // @param in, the input value
-Value sign(HalContext* ctx, const Value& x);
+Value sign(SPUContext* ctx, const Value& x);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/prot_wrapper.cc b/libspu/kernel/hal/prot_wrapper.cc
index b6f11ad2..c367a95a 100644
--- a/libspu/kernel/hal/prot_wrapper.cc
+++ b/libspu/kernel/hal/prot_wrapper.cc
@@ -28,19 +28,6 @@
 namespace spu::kernel::hal {
 namespace {
 
-Value unflattenValue(const ArrayRef& arr, absl::Span<const int64_t> shape) {
-  // The underline MPC engine does not take care of dtype information, so we
-  // should set it as INVALID and let the upper layer to handle it.
-  return Value(unflatten(arr, shape), DT_INVALID);
-}
-
-ArrayRef flattenValue(const Value& v) {
-  // TODO: optimization point.
-  // we can set number of valids bits (according to dtype) to inform lower layer
-  // for optimization. Currently dtype information is not used.
-  return flatten(v.data());
-}
-
 std::tuple<int64_t, int64_t, int64_t> deduceMmulArgs(
     const std::vector<int64_t>& lhs, const std::vector<int64_t>& rhs) {
   SPU_ENFORCE(!lhs.empty() && lhs.size() <= 2);
@@ -64,76 +51,74 @@ std::tuple<int64_t, int64_t, int64_t> deduceMmulArgs(
 
 }  // namespace
 
-#define MAP_UNARY_OP(NAME)                               \
-  Value _##NAME(HalContext* ctx, const Value& in) {      \
-    SPU_TRACE_HAL_DISP(ctx, in);                         \
-    auto ret = mpc::NAME(ctx->prot(), flattenValue(in)); \
-    return unflattenValue(ret, in.shape());              \
+#define MAP_UNARY_OP(NAME)                          \
+  Value _##NAME(SPUContext* ctx, const Value& in) { \
+    SPU_TRACE_HAL_DISP(ctx, in);                    \
+    return mpc::NAME(ctx, in);                      \
   }
 
 #define MAP_SHIFT_OP(NAME)                                       \
-  Value _##NAME(HalContext* ctx, const Value& in, size_t bits) { \
+  Value _##NAME(SPUContext* ctx, const Value& in, size_t bits) { \
     SPU_TRACE_HAL_DISP(ctx, in, bits);                           \
-    auto ret = mpc::NAME(ctx->prot(), flattenValue(in), bits);   \
-    return unflattenValue(ret, in.shape());                      \
+    auto ret = mpc::NAME(ctx, in, bits);                         \
+    return ret;                                                  \
   }
 
 #define MAP_BITREV_OP(NAME)                                                   \
-  Value _##NAME(HalContext* ctx, const Value& in, size_t start, size_t end) { \
+  Value _##NAME(SPUContext* ctx, const Value& in, size_t start, size_t end) { \
     SPU_TRACE_HAL_DISP(ctx, in, start, end);                                  \
-    auto ret = mpc::NAME(ctx->prot(), flattenValue(in), start, end);          \
-    return unflattenValue(ret, in.shape());                                   \
+    auto ret = mpc::NAME(ctx, in, start, end);                                \
+    return ret;                                                               \
   }
 
-#define MAP_BINARY_OP(NAME)                                              \
-  Value _##NAME(HalContext* ctx, const Value& x, const Value& y) {       \
-    SPU_TRACE_HAL_DISP(ctx, x, y);                                       \
-    SPU_ENFORCE(x.shape() == y.shape(), "shape mismatch: x={}, y={}",    \
-                x.shape(), y.shape());                                   \
-    auto ret = mpc::NAME(ctx->prot(), flattenValue(x), flattenValue(y)); \
-    return unflattenValue(ret, x.shape());                               \
+#define MAP_BINARY_OP(NAME)                                           \
+  Value _##NAME(SPUContext* ctx, const Value& x, const Value& y) {    \
+    SPU_TRACE_HAL_DISP(ctx, x, y);                                    \
+    SPU_ENFORCE(x.shape() == y.shape(), "shape mismatch: x={}, y={}", \
+                x.shape(), y.shape());                                \
+    auto ret = mpc::NAME(ctx, x, y);                                  \
+    return ret;                                                       \
   }
 
-#define MAP_MMUL_OP(NAME)                                                  \
-  Value _##NAME(HalContext* ctx, const Value& x, const Value& y) {         \
-    SPU_TRACE_HAL_DISP(ctx, x, y);                                         \
-    auto [m, n, k] = deduceMmulArgs(x.shape(), y.shape());                 \
-    auto ret =                                                             \
-        mpc::NAME(ctx->prot(), flattenValue(x), flattenValue(y), m, n, k); \
-    return unflattenValue(ret, {m, n});                                    \
+#define MAP_MMUL_OP(NAME)                                          \
+  Value _##NAME(SPUContext* ctx, const Value& x, const Value& y) { \
+    SPU_TRACE_HAL_DISP(ctx, x, y);                                 \
+    auto [m, n, k] = deduceMmulArgs(x.shape(), y.shape());         \
+    auto ret = mpc::NAME(ctx, x, y, m, n, k);                      \
+    return ret;                                                    \
   }
 
-Type _common_type_s(HalContext* ctx, const Type& a, const Type& b) {
+Type _common_type_s(SPUContext* ctx, const Type& a, const Type& b) {
   SPU_TRACE_HAL_DISP(ctx, a, b);
-  return mpc::common_type_s(ctx->prot(), a, b);
+  return mpc::common_type_s(ctx, a, b);
 }
 
-Value _cast_type_s(HalContext* ctx, const Value& in, const Type& to) {
+Value _cast_type_s(SPUContext* ctx, const Value& in, const Type& to) {
   SPU_TRACE_HAL_DISP(ctx, in, to);
-  auto ret = mpc::cast_type_s(ctx->prot(), flattenValue(in), to);
-  return unflattenValue(ret, in.shape());
+  auto ret = mpc::cast_type_s(ctx, in, to);
+  return ret;
 }
 
-Value _make_p(HalContext* ctx, uint128_t init,
+Value _make_p(SPUContext* ctx, uint128_t init,
               absl::Span<const int64_t> shape) {
   SPU_TRACE_HAL_DISP(ctx, init);
-  auto res = mpc::make_p(ctx->prot(), init, calcNumel(shape));
-  return unflattenValue(res, shape);
+  auto res = mpc::make_p(ctx, init, Shape(shape));
+  return res;
 }
 
-Value _rand_p(HalContext* ctx, absl::Span<const int64_t> shape) {
+Value _rand_p(SPUContext* ctx, absl::Span<const int64_t> shape) {
   SPU_TRACE_HAL_DISP(ctx, shape);
-  auto rnd = mpc::rand_p(ctx->prot(), calcNumel(shape));
-  return unflattenValue(rnd, shape);
+  auto rnd = mpc::rand_p(ctx, Shape(shape));
+  return rnd;
 }
 
-Value _rand_s(HalContext* ctx, absl::Span<const int64_t> shape) {
+Value _rand_s(SPUContext* ctx, absl::Span<const int64_t> shape) {
   SPU_TRACE_HAL_DISP(ctx, shape);
-  auto rnd = mpc::rand_s(ctx->prot(), calcNumel(shape));
-  return unflattenValue(rnd, shape);
+  auto rnd = mpc::rand_s(ctx, Shape(shape));
+  return rnd;
 }
 
-Value _conv2d_ss(HalContext* ctx, Value input, const Value& kernel,
+Value _conv2d_ss(SPUContext* ctx, Value input, const Value& kernel,
                  absl::Span<const int64_t> window_strides,
                  absl::Span<const int64_t> result_shape) {
   SPU_TRACE_HAL_DISP(ctx, input, kernel, window_strides, result_shape);
@@ -171,24 +156,19 @@ Value _conv2d_ss(HalContext* ctx, Value input, const Value& kernel,
   size_t H = input.shape()[1];
   size_t W = input.shape()[2];
   // FIXME(juhou): define conv2d_ss in api.h to capture this
-  return unflattenValue(
-      ctx->prot()->call("conv2d_aa", flattenValue(input), flattenValue(kernel),
-                        N, H, W, C, O, h, w, stride_h, stride_w),
-      result_shape);
+  return dynDispatch(ctx, "conv2d_aa", input, kernel, N, H, W, C, O, h, w,
+                     stride_h, stride_w);
 }
 
-Value _trunc_p_with_sign(HalContext* ctx, const Value& in, size_t bits,
+Value _trunc_p_with_sign(SPUContext* ctx, const Value& in, size_t bits,
                          bool /*dummy*/) {
   return _trunc_p(ctx, in, bits);
 }
 
-Value _trunc_s_with_sign(HalContext* ctx, const Value& in, size_t bits,
+Value _trunc_s_with_sign(SPUContext* ctx, const Value& in, size_t bits,
                          bool is_positive) {
-  if (ctx->rt_config().protocol() == ProtocolKind::CHEETAH) {
-    return unflattenValue(
-        ctx->prot()->call("trunc_a_with_sign", flattenValue(in), bits,
-                          is_positive),
-        in.shape());
+  if (ctx->config().protocol() == ProtocolKind::CHEETAH) {
+    return dynDispatch(ctx, "trunc_a_with_sign", in, bits, is_positive);
   } else {
     return _trunc_s(ctx, in, bits);
   }
@@ -226,17 +206,17 @@ MAP_MMUL_OP(mmul_pp)
 MAP_MMUL_OP(mmul_sp)
 MAP_MMUL_OP(mmul_ss)
 
-#define MAP_OPTIONAL_BINARY_OP(NAME)                                     \
-  std::optional<Value> _##NAME(HalContext* ctx, const Value& x,          \
-                               const Value& y) {                         \
-    SPU_TRACE_HAL_DISP(ctx, x, y);                                       \
-    SPU_ENFORCE(x.shape() == y.shape(), "shape mismatch: x={}, y={}",    \
-                x.shape(), y.shape());                                   \
-    auto ret = mpc::NAME(ctx->prot(), flattenValue(x), flattenValue(y)); \
-    if (!ret.has_value()) {                                              \
-      return std::nullopt;                                               \
-    }                                                                    \
-    return unflattenValue(ret.value(), x.shape());                       \
+#define MAP_OPTIONAL_BINARY_OP(NAME)                                  \
+  std::optional<Value> _##NAME(SPUContext* ctx, const Value& x,       \
+                               const Value& y) {                      \
+    SPU_TRACE_HAL_DISP(ctx, x, y);                                    \
+    SPU_ENFORCE(x.shape() == y.shape(), "shape mismatch: x={}, y={}", \
+                x.shape(), y.shape());                                \
+    auto ret = mpc::NAME(ctx, x, y);                                  \
+    if (!ret.has_value()) {                                           \
+      return std::nullopt;                                            \
+    }                                                                 \
+    return ret.value();                                               \
   }
 
 MAP_OPTIONAL_BINARY_OP(equal_ss)
diff --git a/libspu/kernel/hal/prot_wrapper.h b/libspu/kernel/hal/prot_wrapper.h
index 977889bb..be890cc3 100644
--- a/libspu/kernel/hal/prot_wrapper.h
+++ b/libspu/kernel/hal/prot_wrapper.h
@@ -14,73 +14,73 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hal {
 
 // !!please read [README.md] for api naming conventions.
-Type _common_type_s(HalContext* ctx, const Type& a, const Type& b);
-Value _cast_type_s(HalContext* ctx, const Value& in, const Type& to);
+Type _common_type_s(SPUContext* ctx, const Type& a, const Type& b);
+Value _cast_type_s(SPUContext* ctx, const Value& in, const Type& to);
 
-Value _p2s(HalContext* ctx, const Value& in);
-Value _s2p(HalContext* ctx, const Value& in);
+Value _p2s(SPUContext* ctx, const Value& in);
+Value _s2p(SPUContext* ctx, const Value& in);
 
-Value _not_p(HalContext* ctx, const Value& in);
-Value _not_s(HalContext* ctx, const Value& in);
+Value _not_p(SPUContext* ctx, const Value& in);
+Value _not_s(SPUContext* ctx, const Value& in);
 
-Value _msb_p(HalContext* ctx, const Value& in);
-Value _msb_s(HalContext* ctx, const Value& in);
+Value _msb_p(SPUContext* ctx, const Value& in);
+Value _msb_s(SPUContext* ctx, const Value& in);
 
-Value _equal_pp(HalContext* ctx, const Value& x, const Value& y);
-std::optional<Value> _equal_sp(HalContext* ctx, const Value& x, const Value& y);
-std::optional<Value> _equal_ss(HalContext* ctx, const Value& x, const Value& y);
+Value _equal_pp(SPUContext* ctx, const Value& x, const Value& y);
+std::optional<Value> _equal_sp(SPUContext* ctx, const Value& x, const Value& y);
+std::optional<Value> _equal_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _lshift_p(HalContext* ctx, const Value& in, size_t bits);
-Value _lshift_s(HalContext* ctx, const Value& in, size_t bits);
+Value _lshift_p(SPUContext* ctx, const Value& in, size_t bits);
+Value _lshift_s(SPUContext* ctx, const Value& in, size_t bits);
 
-Value _rshift_p(HalContext* ctx, const Value& in, size_t bits);
-Value _rshift_s(HalContext* ctx, const Value& in, size_t bits);
+Value _rshift_p(SPUContext* ctx, const Value& in, size_t bits);
+Value _rshift_s(SPUContext* ctx, const Value& in, size_t bits);
 
-Value _arshift_p(HalContext* ctx, const Value& in, size_t bits);
-Value _arshift_s(HalContext* ctx, const Value& in, size_t bits);
-Value _trunc_p(HalContext* ctx, const Value& in, size_t bits);
-Value _trunc_s(HalContext* ctx, const Value& in, size_t bits);
-Value _trunc_p_with_sign(HalContext* ctx, const Value& in, size_t bits,
+Value _arshift_p(SPUContext* ctx, const Value& in, size_t bits);
+Value _arshift_s(SPUContext* ctx, const Value& in, size_t bits);
+Value _trunc_p(SPUContext* ctx, const Value& in, size_t bits);
+Value _trunc_s(SPUContext* ctx, const Value& in, size_t bits);
+Value _trunc_p_with_sign(SPUContext* ctx, const Value& in, size_t bits,
                          bool is_positive);
-Value _trunc_s_with_sign(HalContext* ctx, const Value& in, size_t bits,
+Value _trunc_s_with_sign(SPUContext* ctx, const Value& in, size_t bits,
                          bool is_positive);
 
-Value _add_pp(HalContext* ctx, const Value& x, const Value& y);
-Value _add_sp(HalContext* ctx, const Value& x, const Value& y);
-Value _add_ss(HalContext* ctx, const Value& x, const Value& y);
+Value _add_pp(SPUContext* ctx, const Value& x, const Value& y);
+Value _add_sp(SPUContext* ctx, const Value& x, const Value& y);
+Value _add_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _mul_pp(HalContext* ctx, const Value& x, const Value& y);
-Value _mul_sp(HalContext* ctx, const Value& x, const Value& y);
-Value _mul_ss(HalContext* ctx, const Value& x, const Value& y);
+Value _mul_pp(SPUContext* ctx, const Value& x, const Value& y);
+Value _mul_sp(SPUContext* ctx, const Value& x, const Value& y);
+Value _mul_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _mmul_pp(HalContext* ctx, const Value& x, const Value& y);
-Value _mmul_sp(HalContext* ctx, const Value& x, const Value& y);
-Value _mmul_ss(HalContext* ctx, const Value& x, const Value& y);
+Value _mmul_pp(SPUContext* ctx, const Value& x, const Value& y);
+Value _mmul_sp(SPUContext* ctx, const Value& x, const Value& y);
+Value _mmul_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _conv2d_ss(HalContext* ctx, Value x, const Value& y,
+Value _conv2d_ss(SPUContext* ctx, Value x, const Value& y,
                  absl::Span<const int64_t> window_strides,
                  absl::Span<const int64_t> result_shape);
 
-Value _and_pp(HalContext* ctx, const Value& x, const Value& y);
-Value _and_sp(HalContext* ctx, const Value& x, const Value& y);
-Value _and_ss(HalContext* ctx, const Value& x, const Value& y);
+Value _and_pp(SPUContext* ctx, const Value& x, const Value& y);
+Value _and_sp(SPUContext* ctx, const Value& x, const Value& y);
+Value _and_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _xor_pp(HalContext* ctx, const Value& x, const Value& y);
-Value _xor_sp(HalContext* ctx, const Value& x, const Value& y);
-Value _xor_ss(HalContext* ctx, const Value& x, const Value& y);
+Value _xor_pp(SPUContext* ctx, const Value& x, const Value& y);
+Value _xor_sp(SPUContext* ctx, const Value& x, const Value& y);
+Value _xor_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _bitrev_p(HalContext* ctx, const Value& in, size_t start, size_t end);
-Value _bitrev_s(HalContext* ctx, const Value& in, size_t start, size_t end);
+Value _bitrev_p(SPUContext* ctx, const Value& in, size_t start, size_t end);
+Value _bitrev_s(SPUContext* ctx, const Value& in, size_t start, size_t end);
 
-Value _make_p(HalContext* ctx, uint128_t init, absl::Span<const int64_t> shape);
+Value _make_p(SPUContext* ctx, uint128_t init, absl::Span<const int64_t> shape);
 
-Value _rand_p(HalContext* ctx, absl::Span<const int64_t> shape);
-Value _rand_s(HalContext* ctx, absl::Span<const int64_t> shape);
+Value _rand_p(SPUContext* ctx, absl::Span<const int64_t> shape);
+Value _rand_s(SPUContext* ctx, absl::Span<const int64_t> shape);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/public_helper.cc b/libspu/kernel/hal/public_helper.cc
index 2d134cb9..36096972 100644
--- a/libspu/kernel/hal/public_helper.cc
+++ b/libspu/kernel/hal/public_helper.cc
@@ -20,7 +20,7 @@
 
 namespace spu::kernel::hal {
 
-NdArrayRef dump_public(HalContext* ctx, const Value& v) {
+NdArrayRef dump_public(SPUContext* ctx, const Value& v) {
   SPU_TRACE_HAL_DISP(ctx, v);
   SPU_ENFORCE(v.storage_type().isa<mpc::Pub2kTy>(), "got {}", v.storage_type());
   const auto field = v.storage_type().as<Ring2k>()->field();
diff --git a/libspu/kernel/hal/public_helper.h b/libspu/kernel/hal/public_helper.h
index 4201ea78..01e5d9ad 100644
--- a/libspu/kernel/hal/public_helper.h
+++ b/libspu/kernel/hal/public_helper.h
@@ -14,18 +14,18 @@
 
 #pragma once
 
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 #include "libspu/core/xt_helper.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
 
 namespace spu::kernel::hal {
 
 // Export a value to a buffer.
-NdArrayRef dump_public(HalContext* ctx, const Value& v);
+NdArrayRef dump_public(SPUContext* ctx, const Value& v);
 
 // TODO: Find a more proper space.
 template <typename T>
-xt::xarray<T> dump_public_as(HalContext* ctx, const Value& in) {
+xt::xarray<T> dump_public_as(SPUContext* ctx, const Value& in) {
   auto arr = dump_public(ctx, in);
 
 #define CASE(NAME, TYPE, _)                  \
@@ -44,7 +44,7 @@ xt::xarray<T> dump_public_as(HalContext* ctx, const Value& in) {
 }
 
 template <typename T>
-T getScalarValue(HalContext* ctx, const spu::Value& value) {
+T getScalarValue(SPUContext* ctx, const spu::Value& value) {
   SPU_ENFORCE(value.numel() == 1, "{} is not a scalar tensor.", value);
   SPU_ENFORCE(value.isPublic(), "{} is not a public value", value);
 
diff --git a/libspu/kernel/hal/random.cc b/libspu/kernel/hal/random.cc
index 21e07879..56ad1f48 100644
--- a/libspu/kernel/hal/random.cc
+++ b/libspu/kernel/hal/random.cc
@@ -26,8 +26,8 @@
 namespace spu::kernel::hal {
 namespace {
 
-uint64_t genPublicRandSeed(HalContext* hctx) {
-  auto* prg_state = hctx->prot()->getState<mpc::PrgState>();
+uint64_t genPublicRandSeed(SPUContext* sctx) {
+  auto* prg_state = sctx->prot()->getState<mpc::PrgState>();
   uint64_t seed;
   prg_state->fillPubl(absl::MakeSpan(&seed, 1));
   return seed;
@@ -35,7 +35,7 @@ uint64_t genPublicRandSeed(HalContext* hctx) {
 
 }  // namespace
 
-Value rng_uniform(HalContext* ctx, const Value& lo, const Value& hi,
+Value rng_uniform(SPUContext* ctx, const Value& lo, const Value& hi,
                   absl::Span<const int64_t> to_shape) {
   SPU_TRACE_HAL_LEAF(ctx, lo, hi, to_shape);
   SPU_ENFORCE(lo.isPublic() && hi.isPublic());
@@ -56,7 +56,7 @@ Value rng_uniform(HalContext* ctx, const Value& lo, const Value& hi,
   return constant(ctx, buffer, DT_FXP);
 }
 
-Value random(HalContext* ctx, Visibility vis, DataType dtype,
+Value random(SPUContext* ctx, Visibility vis, DataType dtype,
              absl::Span<const int64_t> shape) {
   Value ret;
   if (vis == VIS_PUBLIC) {
diff --git a/libspu/kernel/hal/random.h b/libspu/kernel/hal/random.h
index b2415bcc..259f60e4 100644
--- a/libspu/kernel/hal/random.h
+++ b/libspu/kernel/hal/random.h
@@ -14,8 +14,8 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hal {
 
@@ -23,13 +23,13 @@ namespace spu::kernel::hal {
 // @param lo, lower limit (include)
 // @param hi, upper limit (exclude)
 // @param to_shape, the target shape
-Value rng_uniform(HalContext* ctx, const Value& lo, const Value& hi,
+Value rng_uniform(SPUContext* ctx, const Value& lo, const Value& hi,
                   absl::Span<const int64_t> to_shape);
 
 /// Make a random value.
 //
 // The value is uniformly distributed in value's range.
-Value random(HalContext* ctx, Visibility vis, DataType dtype,
+Value random(SPUContext* ctx, Visibility vis, DataType dtype,
              absl::Span<const int64_t> shape);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/ring.cc b/libspu/kernel/hal/ring.cc
index fed0e389..0fc2b97b 100644
--- a/libspu/kernel/hal/ring.cc
+++ b/libspu/kernel/hal/ring.cc
@@ -75,12 +75,12 @@ std::tuple<int64_t, int64_t, int64_t> calcMmulTilingSize(int64_t m, int64_t n,
 // shape_ops -> type_cast : concat/pad requires _common_type
 // ring -> shape_ops      : tiled_mmul(transpose/slice)
 // type_cast -> ring      : _p2s/arshift/shift
-Value transpose(HalContext* ctx, const Value& in) {
+Value transpose(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
   return Value(in.data().transpose({}), in.dtype());
 }
 
-Value slice(HalContext* ctx, const Value& in,
+Value slice(SPUContext* ctx, const Value& in,
             absl::Span<const int64_t> start_indices,
             absl::Span<const int64_t> end_indices,
             absl::Span<const int64_t> strides) {
@@ -92,7 +92,7 @@ Value slice(HalContext* ctx, const Value& in,
 
 }  // namespace
 
-Type _common_type(HalContext* ctx, const Type& a, const Type& b) {
+Type _common_type(SPUContext* ctx, const Type& a, const Type& b) {
   if (a.isa<Secret>() && b.isa<Secret>()) {
     return _common_type_s(ctx, a, b);
   } else if (a.isa<Secret>()) {
@@ -105,7 +105,7 @@ Type _common_type(HalContext* ctx, const Type& a, const Type& b) {
   }
 }
 
-Value _cast_type(HalContext* ctx, const Value& x, const Type& to) {
+Value _cast_type(SPUContext* ctx, const Value& x, const Type& to) {
   if (x.isPublic() && to.isa<Public>()) {
     return x;
   } else if (x.isPublic() && to.isa<Secret>()) {
@@ -119,7 +119,7 @@ Value _cast_type(HalContext* ctx, const Value& x, const Type& to) {
 }
 
 #define IMPL_UNARY_OP(Name, FnP, FnS)                       \
-  Value Name(HalContext* ctx, const Value& in) {            \
+  Value Name(SPUContext* ctx, const Value& in) {            \
     SPU_TRACE_HAL_LEAF(ctx, in);                            \
     if (in.isPublic()) {                                    \
       return FnP(ctx, in);                                  \
@@ -131,7 +131,7 @@ Value _cast_type(HalContext* ctx, const Value& x, const Type& to) {
   }
 
 #define IMPL_SHIFT_OP(Name, FnP, FnS)                         \
-  Value Name(HalContext* ctx, const Value& in, size_t bits) { \
+  Value Name(SPUContext* ctx, const Value& in, size_t bits) { \
     SPU_TRACE_HAL_LEAF(ctx, in, bits);                        \
     if (in.isPublic()) {                                      \
       return FnP(ctx, in, bits);                              \
@@ -143,7 +143,7 @@ Value _cast_type(HalContext* ctx, const Value& x, const Type& to) {
   }
 
 #define IMPL_COMMUTATIVE_BINARY_OP(Name, FnPP, FnSP, FnSS)        \
-  Value Name(HalContext* ctx, const Value& x, const Value& y) {   \
+  Value Name(SPUContext* ctx, const Value& x, const Value& y) {   \
     SPU_TRACE_HAL_LEAF(ctx, x, y);                                \
     if (x.isPublic() && y.isPublic()) {                           \
       return FnPP(ctx, x, y);                                     \
@@ -171,13 +171,13 @@ IMPL_COMMUTATIVE_BINARY_OP(_mul, _mul_pp, _mul_sp, _mul_ss)
 IMPL_COMMUTATIVE_BINARY_OP(_and, _and_pp, _and_sp, _and_ss)
 IMPL_COMMUTATIVE_BINARY_OP(_xor, _xor_pp, _xor_sp, _xor_ss)
 
-Value _sub(HalContext* ctx, const Value& x, const Value& y) {
+Value _sub(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
   return _add(ctx, x, _negate(ctx, y));
 }
 
 // TODO: remove this kernel, the algorithm could be used for boolean equal test.
-[[maybe_unused]] Value _eqz(HalContext* ctx, const Value& x) {
+[[maybe_unused]] Value _eqz(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   // eqz(x) = not(lsb(pre_or(x)))
@@ -193,7 +193,7 @@ Value _sub(HalContext* ctx, const Value& x, const Value& y) {
   return res;
 }
 
-Value _conv2d(HalContext* ctx, const Value& x, const Value& y,
+Value _conv2d(SPUContext* ctx, const Value& x, const Value& y,
               absl::Span<const int64_t> window_strides,
               absl::Span<const int64_t> result_shape) {
   // s*p and p*p should call `dot`
@@ -201,7 +201,7 @@ Value _conv2d(HalContext* ctx, const Value& x, const Value& y,
   return _conv2d_ss(ctx, x, y, window_strides, result_shape);
 }
 
-static Value _mmul_impl(HalContext* ctx, const Value& x, const Value& y) {
+static Value _mmul_impl(SPUContext* ctx, const Value& x, const Value& y) {
   if (x.isPublic() && y.isPublic()) {
     return _mmul_pp(ctx, x, y);
   } else if (x.isSecret() && y.isPublic()) {
@@ -215,12 +215,12 @@ static Value _mmul_impl(HalContext* ctx, const Value& x, const Value& y) {
   }
 };
 
-Value _mmul(HalContext* ctx, const Value& x, const Value& y) {
+Value _mmul(SPUContext* ctx, const Value& x, const Value& y) {
   auto [m, n, k] = deduceMmulArgs(x.shape(), y.shape());
   auto [m_step, n_step, k_step] =
       calcMmulTilingSize(m, n, k, x.elsize(), 256UL * 1024 * 1024);
 
-  if (ctx->rt_config().experimental_disable_mmul_split() ||
+  if (ctx->config().experimental_disable_mmul_split() ||
       (m_step == m && n_step == n && k_step == k)) {
     // no split
     return _mmul_impl(ctx, x, y);
@@ -302,12 +302,12 @@ Value _mmul(HalContext* ctx, const Value& x, const Value& y) {
   return ret;
 }
 
-Value _or(HalContext* ctx, const Value& x, const Value& y) {
+Value _or(SPUContext* ctx, const Value& x, const Value& y) {
   // X or Y = X xor Y xor (X and Y)
   return _xor(ctx, x, _xor(ctx, y, _and(ctx, x, y)));
 }
 
-static std::optional<Value> _equal_impl(HalContext* ctx, const Value& x,
+static std::optional<Value> _equal_impl(SPUContext* ctx, const Value& x,
                                         const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
@@ -324,7 +324,7 @@ static std::optional<Value> _equal_impl(HalContext* ctx, const Value& x,
   return std::nullopt;
 }
 
-Value _equal(HalContext* ctx, const Value& x, const Value& y) {
+Value _equal(SPUContext* ctx, const Value& x, const Value& y) {
   // First try use equal kernel, i.e. for 2PC , equal can be done with the same
   // cost of half MSB.
   //      x0 + x1 = y0 + y1 mod 2^k
@@ -343,7 +343,7 @@ Value _equal(HalContext* ctx, const Value& x, const Value& y) {
               _xor(ctx, _less(ctx, y, x), _k1));
 }
 
-Value _trunc_with_sign(HalContext* ctx, const Value& x, size_t bits,
+Value _trunc_with_sign(SPUContext* ctx, const Value& x, size_t bits,
                        bool is_positive) {
   SPU_TRACE_HAL_LEAF(ctx, x, bits);
   bits = (bits == 0) ? ctx->getFxpBits() : bits;
@@ -357,7 +357,7 @@ Value _trunc_with_sign(HalContext* ctx, const Value& x, size_t bits,
   }
 }
 
-Value _trunc(HalContext* ctx, const Value& x, size_t bits) {
+Value _trunc(SPUContext* ctx, const Value& x, size_t bits) {
   SPU_TRACE_HAL_LEAF(ctx, x, bits);
   bits = (bits == 0) ? ctx->getFxpBits() : bits;
 
@@ -370,14 +370,14 @@ Value _trunc(HalContext* ctx, const Value& x, size_t bits) {
   }
 }
 
-Value _negate(HalContext* ctx, const Value& x) {
+Value _negate(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   // negate(x) = not(x) + 1
   return _add(ctx, _not(ctx, x), _constant(ctx, 1, x.shape()));
 }
 
-Value _sign(HalContext* ctx, const Value& x) {
+Value _sign(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   // is_negative = x < 0 ? 1 : 0;
@@ -393,7 +393,7 @@ Value _sign(HalContext* ctx, const Value& x) {
   return _sub(ctx, one, _mul(ctx, two, is_negative));
 }
 
-Value _less(HalContext* ctx, const Value& x, const Value& y) {
+Value _less(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
   // Note: the impl assume inputs are signed with two's complement encoding.
@@ -402,7 +402,7 @@ Value _less(HalContext* ctx, const Value& x, const Value& y) {
 }
 
 // swap bits of [start, end)
-Value _bitrev(HalContext* ctx, const Value& x, size_t start, size_t end) {
+Value _bitrev(SPUContext* ctx, const Value& x, size_t start, size_t end) {
   SPU_TRACE_HAL_LEAF(ctx, x, start, end);
 
   if (x.isPublic()) {
@@ -414,14 +414,14 @@ Value _bitrev(HalContext* ctx, const Value& x, size_t start, size_t end) {
   SPU_THROW("unsupport op={} for {}", "_bitrev", x);
 }
 
-Value _mux(HalContext* ctx, const Value& pred, const Value& a, const Value& b) {
+Value _mux(SPUContext* ctx, const Value& pred, const Value& a, const Value& b) {
   SPU_TRACE_HAL_LEAF(ctx, pred, a, b);
 
   // b + pred*(a-b)
   return _add(ctx, b, _mul(ctx, pred, _sub(ctx, a, b)));
 }
 
-Value _clamp(HalContext* ctx, const Value& x, const Value& minv,
+Value _clamp(SPUContext* ctx, const Value& x, const Value& minv,
              const Value& maxv) {
   SPU_TRACE_HAL_LEAF(ctx, x, minv, maxv);
 
@@ -432,12 +432,12 @@ Value _clamp(HalContext* ctx, const Value& x, const Value& minv,
   return _mux(ctx, _less(ctx, res, maxv), res, maxv);
 }
 
-Value _constant(HalContext* ctx, uint128_t init,
+Value _constant(SPUContext* ctx, uint128_t init,
                 absl::Span<const int64_t> shape) {
   return _make_p(ctx, init, shape);
 }
 
-Value _bit_parity(HalContext* ctx, const Value& x, size_t bits) {
+Value _bit_parity(SPUContext* ctx, const Value& x, size_t bits) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(absl::has_single_bit(bits), "currently only support power of 2");
@@ -452,7 +452,7 @@ Value _bit_parity(HalContext* ctx, const Value& x, size_t bits) {
 }
 
 // TODO(jint): OPTIMIZE ME, this impl seems to be super slow.
-Value _popcount(HalContext* ctx, const Value& x, size_t bits) {
+Value _popcount(SPUContext* ctx, const Value& x, size_t bits) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   Value ret = _constant(ctx, 0, x.shape());
@@ -478,7 +478,7 @@ Value _popcount(HalContext* ctx, const Value& x, size_t bits) {
 //   x2:  011110000   ; x1 | (x1>>2)
 //   x3:  011111111   ; x2 | (x2>>4)
 //
-Value _prefix_or(HalContext* ctx, const Value& x) {
+Value _prefix_or(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   auto b0 = _prefer_b(ctx, x);
@@ -491,7 +491,7 @@ Value _prefix_or(HalContext* ctx, const Value& x) {
   return b0;
 }
 
-Value _bitdeintl(HalContext* ctx, const Value& in) {
+Value _bitdeintl(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_LEAF(ctx, in);
 
   // algorithm:
@@ -530,9 +530,9 @@ Value _bitdeintl(HalContext* ctx, const Value& in) {
 // - b2a+2*mul > 2*mula2b, we prefer to leave p as bshare.
 //
 // Cheetah is the later case.
-Value _prefer_a(HalContext* ctx, const Value& x) {
+Value _prefer_a(SPUContext* ctx, const Value& x) {
   if (x.storage_type().isa<BShare>()) {
-    if (ctx->rt_config().protocol() == ProtocolKind::CHEETAH &&
+    if (ctx->config().protocol() == ProtocolKind::CHEETAH &&
         x.storage_type().as<BShare>()->nbits() == 1) {
       return x;
     }
@@ -550,7 +550,7 @@ Value _prefer_a(HalContext* ctx, const Value& x) {
 //   y2 = v ^ c2
 // When a variable is followed by multiple binary operation, it's more efficient
 // to convert it to boolean share first.
-Value _prefer_b(HalContext* ctx, const Value& x) {
+Value _prefer_b(SPUContext* ctx, const Value& x) {
   if (x.storage_type().isa<AShare>()) {
     const auto k0 = _constant(ctx, 0U, x.shape());
     return _xor(ctx, x, k0).setDtype(x.dtype());  // noop, to bshare
diff --git a/libspu/kernel/hal/ring.h b/libspu/kernel/hal/ring.h
index ae010bc3..a7e2d2d5 100644
--- a/libspu/kernel/hal/ring.h
+++ b/libspu/kernel/hal/ring.h
@@ -14,8 +14,8 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 // !!please read [README.md] for api naming conventions.
 // this module implements ops x ring 2k space WITHOUT dtype check.
@@ -26,92 +26,92 @@
 
 namespace spu::kernel::hal {
 
-Type _common_type(HalContext* ctx, const Type& a, const Type& b);
+Type _common_type(SPUContext* ctx, const Type& a, const Type& b);
 
-Value _cast_type(HalContext* ctx, const Value& x, const Type& to);
+Value _cast_type(SPUContext* ctx, const Value& x, const Type& to);
 
-Value _negate(HalContext* ctx, const Value& x);
+Value _negate(SPUContext* ctx, const Value& x);
 
 // Return 1 when x >= 0 else -1.
-Value _sign(HalContext* ctx, const Value& x);
+Value _sign(SPUContext* ctx, const Value& x);
 
-Value _add(HalContext* ctx, const Value& x, const Value& y);
+Value _add(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _sub(HalContext* ctx, const Value& x, const Value& y);
+Value _sub(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _mul(HalContext* ctx, const Value& x, const Value& y);
+Value _mul(SPUContext* ctx, const Value& x, const Value& y);
 
 // Note: there is no div (multiplicative inverse) in ring 2^k.
-// Value _div(HalContext* ctx, const Value& x, const Value& y);
+// Value _div(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _mmul(HalContext* ctx, const Value& x, const Value& y);
+Value _mmul(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _conv2d(HalContext* ctx, const Value& x, const Value& y,
+Value _conv2d(SPUContext* ctx, const Value& x, const Value& y,
               absl::Span<const int64_t> window_strides,
               absl::Span<const int64_t> result_shape);
 
-Value _and(HalContext* ctx, const Value& x, const Value& y);
+Value _and(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _xor(HalContext* ctx, const Value& x, const Value& y);
+Value _xor(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _or(HalContext* ctx, const Value& x, const Value& y);
+Value _or(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _not(HalContext* ctx, const Value& in);
+Value _not(SPUContext* ctx, const Value& in);
 
-Value _msb(HalContext* ctx, const Value& in);
+Value _msb(SPUContext* ctx, const Value& in);
 
 // Return 1{x == y}
-Value _equal(HalContext* ctx, const Value& x, const Value& y);
+Value _equal(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _less(HalContext* ctx, const Value& x, const Value& y);
+Value _less(SPUContext* ctx, const Value& x, const Value& y);
 
-Value _lshift(HalContext* ctx, const Value& in, size_t bits);
+Value _lshift(SPUContext* ctx, const Value& in, size_t bits);
 
-Value _rshift(HalContext* ctx, const Value& in, size_t bits);
+Value _rshift(SPUContext* ctx, const Value& in, size_t bits);
 
-Value _arshift(HalContext* ctx, const Value& in, size_t bits);
+Value _arshift(SPUContext* ctx, const Value& in, size_t bits);
 
-Value _trunc(HalContext* ctx, const Value& x, size_t bits = 0);
+Value _trunc(SPUContext* ctx, const Value& x, size_t bits = 0);
 
-Value _trunc_with_sign(HalContext* ctx, const Value& x, size_t bits = 0,
+Value _trunc_with_sign(SPUContext* ctx, const Value& x, size_t bits = 0,
                        bool is_positive = true);
 
-Value _bitrev(HalContext* ctx, const Value&, size_t start_idx, size_t end_idx);
+Value _bitrev(SPUContext* ctx, const Value&, size_t start_idx, size_t end_idx);
 
 // Expect pred is either {0, 1}.
-Value _mux(HalContext* ctx, const Value& pred, const Value& a, const Value& b);
+Value _mux(SPUContext* ctx, const Value& pred, const Value& a, const Value& b);
 
 // TODO: test me
-Value _clamp(HalContext* ctx, const Value& x, const Value& minv,
+Value _clamp(SPUContext* ctx, const Value& x, const Value& minv,
              const Value& maxv);
 // Make a public value from uint128_t init value.
 //
 // If the current working field has less than 128bit, the lower sizeof(field)
 // bits are used.
-Value _constant(HalContext* ctx, uint128_t init,
+Value _constant(SPUContext* ctx, uint128_t init,
                 absl::Span<const int64_t> shape);
 
 // Return the parity of bits, that is
 // - 1 if there are odd number of 1s.
 // - 0 if there are even number of 1s.
-Value _bit_parity(HalContext* ctx, const Value& x, size_t bits);
+Value _bit_parity(SPUContext* ctx, const Value& x, size_t bits);
 
-Value _popcount(HalContext* ctx, const Value& x, size_t bits);
+Value _popcount(SPUContext* ctx, const Value& x, size_t bits);
 
 // out[i] = OR(in[0..i])
-Value _prefix_or(HalContext* ctx, const Value& x);
+Value _prefix_or(SPUContext* ctx, const Value& x);
 
 // separate even and odd bits. e.g.
 //   xAyBzCwD -> xyzwABCD
-Value _bitdeintl(HalContext* ctx, const Value& in);
+Value _bitdeintl(SPUContext* ctx, const Value& in);
 
 // Return value in arithmetic shared form if it's not.
 //
 // Note: theoretically, we should not leak `share` concept to hal layer.
-Value _prefer_a(HalContext* ctx, const Value& x);
+Value _prefer_a(SPUContext* ctx, const Value& x);
 
 // Return value in binary shared form if it's not.
 // Note: theoretically, we should not leak `share` concept to hal layer.
-Value _prefer_b(HalContext* ctx, const Value& x);
+Value _prefer_b(SPUContext* ctx, const Value& x);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/shape_ops.cc b/libspu/kernel/hal/shape_ops.cc
index 6e1dbf8a..e5e7edae 100644
--- a/libspu/kernel/hal/shape_ops.cc
+++ b/libspu/kernel/hal/shape_ops.cc
@@ -29,7 +29,7 @@
 
 namespace spu::kernel::hal {
 
-Value transpose(HalContext* ctx, const Value& in,
+Value transpose(SPUContext* ctx, const Value& in,
                 absl::Span<const int64_t> permutation) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
@@ -39,7 +39,7 @@ Value transpose(HalContext* ctx, const Value& in,
   return Value(in.data().transpose(permutation), in.dtype()).clone();
 }
 
-Value slice(HalContext* ctx, const Value& in,
+Value slice(SPUContext* ctx, const Value& in,
             absl::Span<const int64_t> start_indices,
             absl::Span<const int64_t> end_indices,
             absl::Span<const int64_t> strides) {
@@ -49,12 +49,12 @@ Value slice(HalContext* ctx, const Value& in,
                in.dtype());
 }
 
-Value slice_scalar_at(HalContext* ctx, const Value& input,
+Value slice_scalar_at(SPUContext* ctx, const Value& input,
                       absl::Span<const int64_t> indices) {
   return Value(input.data().slice_scalar_at(indices), input.dtype());
 }
 
-Value update_slice(HalContext* ctx, const Value& in, const Value& update,
+Value update_slice(SPUContext* ctx, const Value& in, const Value& update,
                    absl::Span<const int64_t> start_indices) {
   auto ret = in.clone();
   auto u = stype_cast(ctx, update, ret.storage_type());
@@ -62,14 +62,14 @@ Value update_slice(HalContext* ctx, const Value& in, const Value& update,
   return ret;
 }
 
-Value reshape(HalContext* ctx, const Value& in,
+Value reshape(SPUContext* ctx, const Value& in,
               absl::Span<const int64_t> to_shape) {
   SPU_TRACE_HAL_DISP(ctx, in, to_shape);
 
   return Value(in.data().reshape(to_shape), in.dtype());
 }
 
-Value broadcast_to(HalContext* ctx, const Value& in,
+Value broadcast_to(SPUContext* ctx, const Value& in,
                    absl::Span<const int64_t> to_shape,
                    absl::Span<const int64_t> in_dims) {
   SPU_TRACE_HAL_DISP(ctx, in, to_shape);
@@ -77,19 +77,19 @@ Value broadcast_to(HalContext* ctx, const Value& in,
   return Value(in.data().broadcast_to(to_shape, in_dims), in.dtype());
 }
 
-Value reverse(HalContext* ctx, const Value& in,
+Value reverse(SPUContext* ctx, const Value& in,
               absl::Span<const int64_t> dimensions) {
   SPU_TRACE_HAL_DISP(ctx, in, dimensions);
 
   return Value(in.data().reverse(dimensions), in.dtype());
 }
 
-Value expand(HalContext* ctx, const Value& in,
+Value expand(SPUContext* ctx, const Value& in,
              absl::Span<const int64_t> to_shape) {
   return Value(in.data().expand(to_shape), in.dtype());
 }
 
-Value pad(HalContext* ctx, const Value& in, const Value& padding_value,
+Value pad(SPUContext* ctx, const Value& in, const Value& padding_value,
           absl::Span<const int64_t> edge_padding_low,
           absl::Span<const int64_t> edge_padding_high,
           absl::Span<const int64_t> interior_padding) {
@@ -108,7 +108,7 @@ Value pad(HalContext* ctx, const Value& in, const Value& padding_value,
                in.dtype());
 }
 
-Value concatenate(HalContext* ctx, absl::Span<const Value> values,
+Value concatenate(SPUContext* ctx, absl::Span<const Value> values,
                   const size_t& axis) {
   SPU_TRACE_HAL_DISP(ctx, axis);
   SPU_ENFORCE(!values.empty(), "got={}", values.size());
diff --git a/libspu/kernel/hal/shape_ops.h b/libspu/kernel/hal/shape_ops.h
index 98aaed15..5d919ab2 100644
--- a/libspu/kernel/hal/shape_ops.h
+++ b/libspu/kernel/hal/shape_ops.h
@@ -16,22 +16,22 @@
 
 #include <cstdint>
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hal {
 
 /// the broadcast function
 // @param in, the input
 // @param to_shape, the target shape
-Value broadcast_to(HalContext* ctx, const Value& in,
+Value broadcast_to(SPUContext* ctx, const Value& in,
                    absl::Span<const int64_t> to_shape,
                    absl::Span<const int64_t> in_dims = {});
 
 /// the reshape function
 // @param in, the input
 // @param to_shape, the target shape
-Value reshape(HalContext* ctx, const Value& in,
+Value reshape(SPUContext* ctx, const Value& in,
               absl::Span<const int64_t> to_shape);
 
 /// the slice function
@@ -39,34 +39,34 @@ Value reshape(HalContext* ctx, const Value& in,
 // @param start_indices, the start indices
 // @param end_indices, the end indices
 // @param strides, the strides
-Value slice(HalContext* ctx, const Value& input,
+Value slice(SPUContext* ctx, const Value& input,
             absl::Span<const int64_t> start_indices,
             absl::Span<const int64_t> end_indices,
             absl::Span<const int64_t> strides);
 
 /// This is a special slice for single element at indices
 // @returns a array with empty shape (scalar)
-Value slice_scalar_at(HalContext* ctx, const Value& input,
+Value slice_scalar_at(SPUContext* ctx, const Value& input,
                       absl::Span<const int64_t> indices);
 
 // update a block of in with update, start_indices is postion at in
-Value update_slice(HalContext* ctx, const Value& in, const Value& update,
+Value update_slice(SPUContext* ctx, const Value& in, const Value& update,
                    absl::Span<const int64_t> start_indices);
 
 /// the transpose function
 // @param in, the param
-Value transpose(HalContext* ctx, const Value& in,
+Value transpose(SPUContext* ctx, const Value& in,
                 absl::Span<const int64_t> permutation = {});
 
 //// the reverse function
 // @param in, the param
 // @param dimensions, dimensions to reverse
-Value reverse(HalContext* ctx, const Value& in,
+Value reverse(SPUContext* ctx, const Value& in,
               absl::Span<const int64_t> dimensions);
 
 /// Expand a scalar into to_shape.
 /// Compare with broadcast, expand actually reallocates and assign memory
-Value expand(HalContext* ctx, const Value& in,
+Value expand(SPUContext* ctx, const Value& in,
              absl::Span<const int64_t> to_shape);
 
 //// the pad function
@@ -78,7 +78,7 @@ Value expand(HalContext* ctx, const Value& in,
 //        (next to the highest index) of each dimension
 // @param interior_padding, the amount of padding added between any two elements
 //        in each dimension
-Value pad(HalContext* ctx, const Value& in, const Value& padding_value,
+Value pad(SPUContext* ctx, const Value& in, const Value& padding_value,
           absl::Span<const int64_t> edge_padding_low,
           absl::Span<const int64_t> edge_padding_high,
           absl::Span<const int64_t> interior_padding);
@@ -87,7 +87,7 @@ Value pad(HalContext* ctx, const Value& in, const Value& padding_value,
 // @param first, the first param
 // @param second, the second param
 // @param axis, the axis
-Value concatenate(HalContext* ctx, absl::Span<const Value> values,
+Value concatenate(SPUContext* ctx, absl::Span<const Value> values,
                   const size_t& axis);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/type_cast.cc b/libspu/kernel/hal/type_cast.cc
index 695ff4da..172463d2 100644
--- a/libspu/kernel/hal/type_cast.cc
+++ b/libspu/kernel/hal/type_cast.cc
@@ -22,7 +22,7 @@
 namespace spu::kernel::hal {
 namespace {
 
-Value int2fxp(HalContext* ctx, const Value& x) {
+Value int2fxp(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
   SPU_ENFORCE(x.isInt(), "expect integer, got {}", x.dtype());
 
@@ -44,7 +44,7 @@ Value int2fxp(HalContext* ctx, const Value& x) {
 //   fxp2int(-1.0) = floor(-1+0.999999) = -1
 //   fxp2int(-1.2) = floor(-1.2+0.9999999) = -1
 //
-Value fxp2int(HalContext* ctx, const Value& x, DataType to_type) {
+Value fxp2int(SPUContext* ctx, const Value& x, DataType to_type) {
   SPU_TRACE_HAL_LEAF(ctx, x);
   SPU_ENFORCE(x.dtype() == DataType::DT_FXP);
 
@@ -60,17 +60,17 @@ Value fxp2int(HalContext* ctx, const Value& x, DataType to_type) {
 }  // namespace
 
 // TODO: move seal/reveal into a new header file.
-Value seal(HalContext* ctx, const Value& x) {
+Value seal(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
   return _p2s(ctx, x).setDtype(x.dtype());
 }
 
-Value reveal(HalContext* ctx, const Value& x) {
+Value reveal(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
   return _s2p(ctx, x).setDtype(x.dtype());
 }
 
-Value dtype_cast(HalContext* ctx, const Value& in, DataType to_type) {
+Value dtype_cast(SPUContext* ctx, const Value& in, DataType to_type) {
   SPU_TRACE_HAL_DISP(ctx, in, to_type);
 
   if (to_type == in.dtype()) {
@@ -102,7 +102,7 @@ Value dtype_cast(HalContext* ctx, const Value& in, DataType to_type) {
   SPU_THROW("should not be here");
 }
 
-Value stype_cast(HalContext* ctx, const Value& in, const Type& to) {
+Value stype_cast(SPUContext* ctx, const Value& in, const Type& to) {
   if (in.storage_type() == to) {
     return in;
   }
diff --git a/libspu/kernel/hal/type_cast.h b/libspu/kernel/hal/type_cast.h
index 83a821e7..29a286f3 100644
--- a/libspu/kernel/hal/type_cast.h
+++ b/libspu/kernel/hal/type_cast.h
@@ -14,25 +14,25 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hal {
 
 /// cast dtype
 // @param in, the input value
 // @param to_type, destination dtype
-Value dtype_cast(HalContext* ctx, const Value& in, DataType to_type);
+Value dtype_cast(SPUContext* ctx, const Value& in, DataType to_type);
 
 /// cast public to secret
 // @param in, the input value
-Value seal(HalContext* ctx, const Value& x);
+Value seal(SPUContext* ctx, const Value& x);
 
 /// reveal a secret
 // @param in, the input value
-Value reveal(HalContext* ctx, const Value& x);
+Value reveal(SPUContext* ctx, const Value& x);
 
 /// cast in storage type to dest type
-Value stype_cast(HalContext* ctx, const Value& in, const Type& to);
+Value stype_cast(SPUContext* ctx, const Value& in, const Type& to);
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hlo/BUILD.bazel b/libspu/kernel/hlo/BUILD.bazel
index f9528e92..9168ea75 100644
--- a/libspu/kernel/hlo/BUILD.bazel
+++ b/libspu/kernel/hlo/BUILD.bazel
@@ -52,7 +52,7 @@ spu_cc_test(
     deps = [
         ":basic_binary",
         ":const",
-        ":test_utils",
+        "//libspu/kernel:test_util",
         "//libspu/kernel/hal",
         "//libspu/mpc/utils:simulate",
     ],
@@ -74,7 +74,7 @@ spu_cc_test(
     deps = [
         ":basic_ternary",
         ":const",
-        ":test_utils",
+        "//libspu/kernel:test_util",
         "//libspu/kernel/hal",
         "//libspu/mpc/utils:simulate",
     ],
@@ -96,7 +96,7 @@ spu_cc_test(
     deps = [
         ":basic_unary",
         ":const",
-        ":test_utils",
+        "//libspu/kernel:test_util",
         "//libspu/kernel/hal",
         "//libspu/mpc/utils:simulate",
     ],
@@ -118,7 +118,7 @@ spu_cc_test(
     deps = [
         ":casting",
         ":const",
-        "//libspu/kernel/hal:test_util",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -138,7 +138,7 @@ spu_cc_test(
     srcs = ["const_test.cc"],
     deps = [
         ":const",
-        "//libspu/kernel/hal:test_util",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -184,7 +184,7 @@ spu_cc_test(
     srcs = ["indexing_test.cc"],
     deps = [
         ":indexing",
-        "//libspu/kernel/hal:test_util",
+        "//libspu/kernel:test_util",
     ],
 )
 
@@ -205,7 +205,7 @@ spu_cc_test(
     deps = [
         ":const",
         ":geometrical",
-        ":test_utils",
+        "//libspu/kernel:test_util",
         "//libspu/mpc/utils:simulate",
     ],
 )
@@ -270,8 +270,8 @@ spu_cc_test(
     srcs = ["sort_test.cc"],
     deps = [
         ":sort",
+        "//libspu/kernel:test_util",
         "//libspu/kernel/hal:polymorphic",
-        "//libspu/kernel/hal:test_util",
     ],
 )
 
@@ -307,14 +307,6 @@ spu_cc_test(
     srcs = ["shuffle_test.cc"],
     deps = [
         ":shuffle",
-        "//libspu/kernel/hal:test_util",
-    ],
-)
-
-spu_cc_library(
-    name = "test_utils",
-    hdrs = ["test_utils.h"],
-    deps = [
-        "//libspu:spu_cc_proto",
+        "//libspu/kernel:test_util",
     ],
 )
diff --git a/libspu/kernel/hlo/basic_binary.cc b/libspu/kernel/hlo/basic_binary.cc
index 050ef2cf..edb5ac9c 100644
--- a/libspu/kernel/hlo/basic_binary.cc
+++ b/libspu/kernel/hlo/basic_binary.cc
@@ -21,7 +21,7 @@
 namespace spu::kernel::hlo {
 
 #define SIMPLE_BINARY_KERNEL_DEFN(NAME, HalFcn)           \
-  spu::Value NAME(HalContext *ctx, const spu::Value &lhs, \
+  spu::Value NAME(SPUContext *ctx, const spu::Value &lhs, \
                   const spu::Value &rhs) {                \
     return HalFcn(ctx, lhs, rhs);                         \
   }
@@ -45,7 +45,7 @@ SIMPLE_BINARY_KERNEL_DEFN(GreaterEqual, hal::greater_equal)
 
 #undef SIMPLE_BINARY_KERNEL_DEFN
 
-spu::Value Remainder(HalContext *ctx, const spu::Value &lhs,
+spu::Value Remainder(SPUContext *ctx, const spu::Value &lhs,
                      const spu::Value &rhs) {
   SPU_ENFORCE(lhs.dtype() == rhs.dtype(), "dtype mismatch {} != {}",
               lhs.dtype(), rhs.dtype());
@@ -64,7 +64,7 @@ spu::Value Remainder(HalContext *ctx, const spu::Value &lhs,
   return hal::sub(ctx, lhs, hal::mul(ctx, quotient, rhs));
 }
 
-spu::Value Dot(HalContext *ctx, const spu::Value &lhs, const spu::Value &rhs) {
+spu::Value Dot(SPUContext *ctx, const spu::Value &lhs, const spu::Value &rhs) {
   SPU_ENFORCE(!lhs.shape().empty() && lhs.shape().size() <= 2);
   SPU_ENFORCE(!rhs.shape().empty() && rhs.shape().size() <= 2);
 
diff --git a/libspu/kernel/hlo/basic_binary.h b/libspu/kernel/hlo/basic_binary.h
index 9c2c5556..7386ccb8 100644
--- a/libspu/kernel/hlo/basic_binary.h
+++ b/libspu/kernel/hlo/basic_binary.h
@@ -14,16 +14,16 @@
 
 #pragma once
 
-#include "libspu/kernel/value.h"
+#include "libspu/core/value.h"
 
 namespace spu {
-class HalContext;
+class SPUContext;
 }
 
 namespace spu::kernel::hlo {
 
 #define SIMPLE_BINARY_KERNEL_DECL(NAME)                   \
-  spu::Value NAME(HalContext *ctx, const spu::Value &lhs, \
+  spu::Value NAME(SPUContext *ctx, const spu::Value &lhs, \
                   const spu::Value &rhs);
 
 SIMPLE_BINARY_KERNEL_DECL(Add)
diff --git a/libspu/kernel/hlo/basic_ternary.cc b/libspu/kernel/hlo/basic_ternary.cc
index a205cf5b..333d1722 100644
--- a/libspu/kernel/hlo/basic_ternary.cc
+++ b/libspu/kernel/hlo/basic_ternary.cc
@@ -18,12 +18,12 @@
 
 namespace spu::kernel::hlo {
 
-spu::Value Select(HalContext *ctx, const spu::Value &pred,
+spu::Value Select(SPUContext *ctx, const spu::Value &pred,
                   const spu::Value &on_true, const spu::Value &on_false) {
   return hal::select(ctx, pred, on_true, on_false);
 }
 
-spu::Value Clamp(HalContext *ctx, const spu::Value &operand,
+spu::Value Clamp(SPUContext *ctx, const spu::Value &operand,
                  const spu::Value &min, const spu::Value &max) {
   return hal::clamp(ctx, operand, min, max);
 }
diff --git a/libspu/kernel/hlo/basic_ternary.h b/libspu/kernel/hlo/basic_ternary.h
index e9067760..e6830d26 100644
--- a/libspu/kernel/hlo/basic_ternary.h
+++ b/libspu/kernel/hlo/basic_ternary.h
@@ -14,18 +14,18 @@
 
 #pragma once
 
-#include "libspu/kernel/value.h"
+#include "libspu/core/value.h"
 
 namespace spu {
-class HalContext;
+class SPUContext;
 }
 
 namespace spu::kernel::hlo {
 
-spu::Value Select(HalContext *ctx, const spu::Value &pred,
+spu::Value Select(SPUContext *ctx, const spu::Value &pred,
                   const spu::Value &on_true, const spu::Value &on_false);
 
-spu::Value Clamp(HalContext *ctx, const spu::Value &operand,
+spu::Value Clamp(SPUContext *ctx, const spu::Value &operand,
                  const spu::Value &min, const spu::Value &max);
 
 }  // namespace spu::kernel::hlo
diff --git a/libspu/kernel/hlo/basic_unary.cc b/libspu/kernel/hlo/basic_unary.cc
index efdbff10..c74efe95 100644
--- a/libspu/kernel/hlo/basic_unary.cc
+++ b/libspu/kernel/hlo/basic_unary.cc
@@ -14,16 +14,16 @@
 
 #include "libspu/kernel/hlo/basic_unary.h"
 
-#include "libspu/kernel/context.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hal/constants.h"
 #include "libspu/kernel/hal/polymorphic.h"
 #include "libspu/kernel/hal/type_cast.h"
-#include "libspu/kernel/value.h"
 
 namespace spu::kernel::hlo {
 
 #define SIMPLE_UNARY_KERNEL_DEFN(NAME, HalFcn)             \
-  spu::Value NAME(HalContext *ctx, const spu::Value &in) { \
+  spu::Value NAME(SPUContext *ctx, const spu::Value &in) { \
     return HalFcn(ctx, in);                                \
   }
 
@@ -42,7 +42,7 @@ SIMPLE_UNARY_KERNEL_DEFN(Sqrt, hal::sqrt)
 
 #undef SIMPLE_UNARY_KERNEL_DEFN
 
-spu::Value Expm1(HalContext *ctx, const spu::Value &in) {
+spu::Value Expm1(SPUContext *ctx, const spu::Value &in) {
   // FIXME: By numpy spec, expm1 should have a higher numeric accuracy compare
   // with exp(x) - 1. SPU is not doing so right now, rethink about what we
   // should do here.
@@ -50,7 +50,7 @@ spu::Value Expm1(HalContext *ctx, const spu::Value &in) {
                   hal::constant(ctx, 1.0, DT_FXP, in.shape()));
 }
 
-spu::Value Not(HalContext *ctx, const spu::Value &in) {
+spu::Value Not(SPUContext *ctx, const spu::Value &in) {
   if (in.dtype() == DT_I1) {
     return hal::logical_not(ctx, in);
   } else {
@@ -60,7 +60,7 @@ spu::Value Not(HalContext *ctx, const spu::Value &in) {
   }
 }
 
-spu::Value Sign(HalContext *ctx, const spu::Value &in) {
+spu::Value Sign(SPUContext *ctx, const spu::Value &in) {
   // get the (-1, 1) sign
   auto s = hal::sign(ctx, in);
 
@@ -71,7 +71,7 @@ spu::Value Sign(HalContext *ctx, const spu::Value &in) {
   return hal::dtype_cast(ctx, s, in.dtype());
 }
 
-spu::Value Round_AFZ(HalContext *ctx, const spu::Value &in) {
+spu::Value Round_AFZ(SPUContext *ctx, const spu::Value &in) {
   // select(x < 0, (int)(x-0.5), (int)(x+0.5))
   // -> (float)(int)(x + sign(x) * 0.5)
   SPU_ENFORCE(in.isFxp(), "Round only supports fxp");
diff --git a/libspu/kernel/hlo/basic_unary.h b/libspu/kernel/hlo/basic_unary.h
index d2bd8e4d..67e3188b 100644
--- a/libspu/kernel/hlo/basic_unary.h
+++ b/libspu/kernel/hlo/basic_unary.h
@@ -19,7 +19,7 @@
 namespace spu::kernel::hlo {
 
 #define SIMPLE_UNARY_KERNEL_DECL(NAME) \
-  spu::Value NAME(HalContext *ctx, const spu::Value &in);
+  spu::Value NAME(SPUContext *ctx, const spu::Value &in);
 
 SIMPLE_UNARY_KERNEL_DECL(Reciprocal)
 SIMPLE_UNARY_KERNEL_DECL(Neg)
diff --git a/libspu/kernel/hlo/casting.cc b/libspu/kernel/hlo/casting.cc
index d504ffc6..f6dc195f 100644
--- a/libspu/kernel/hlo/casting.cc
+++ b/libspu/kernel/hlo/casting.cc
@@ -19,7 +19,7 @@
 
 namespace spu::kernel::hlo {
 
-spu::Value Cast(HalContext *ctx, const spu::Value &in, Visibility dst_vtype,
+spu::Value Cast(SPUContext *ctx, const spu::Value &in, Visibility dst_vtype,
                 DataType dst_dtype) {
   spu::Value ret = in;
   if (ret.vtype() != dst_vtype) {
@@ -35,15 +35,15 @@ spu::Value Cast(HalContext *ctx, const spu::Value &in, Visibility dst_vtype,
   return ret;
 }
 
-spu::Value Bitcast(HalContext *ctx, const spu::Value &in, DataType dst_dtype) {
+spu::Value Bitcast(SPUContext *ctx, const spu::Value &in, DataType dst_dtype) {
   return hal::bitcast(ctx, in, dst_dtype);
 }
 
-spu::Value Reveal(HalContext *ctx, const spu::Value &in) {
+spu::Value Reveal(SPUContext *ctx, const spu::Value &in) {
   return hal::reveal(ctx, in);
 }
 
-spu::Value Seal(HalContext *ctx, const spu::Value &in) {
+spu::Value Seal(SPUContext *ctx, const spu::Value &in) {
   return hal::seal(ctx, in);
 }
 
diff --git a/libspu/kernel/hlo/casting.h b/libspu/kernel/hlo/casting.h
index 87181efb..899512ae 100644
--- a/libspu/kernel/hlo/casting.h
+++ b/libspu/kernel/hlo/casting.h
@@ -18,13 +18,13 @@
 
 namespace spu::kernel::hlo {
 
-spu::Value Cast(HalContext *ctx, const spu::Value &in, Visibility dst_vtype,
+spu::Value Cast(SPUContext *ctx, const spu::Value &in, Visibility dst_vtype,
                 DataType dst_dtype);
 
-spu::Value Bitcast(HalContext *ctx, const spu::Value &in, DataType dst_dtype);
+spu::Value Bitcast(SPUContext *ctx, const spu::Value &in, DataType dst_dtype);
 
-spu::Value Reveal(HalContext *ctx, const spu::Value &in);
+spu::Value Reveal(SPUContext *ctx, const spu::Value &in);
 
-spu::Value Seal(HalContext *ctx, const spu::Value &in);
+spu::Value Seal(SPUContext *ctx, const spu::Value &in);
 
 }  // namespace spu::kernel::hlo
diff --git a/libspu/kernel/hlo/const.cc b/libspu/kernel/hlo/const.cc
index fd0ae8ef..6713539f 100644
--- a/libspu/kernel/hlo/const.cc
+++ b/libspu/kernel/hlo/const.cc
@@ -21,7 +21,7 @@
 namespace spu::kernel::hlo {
 
 // TODO: pass DataType as a parameter?
-spu::Value Constant(HalContext *ctx, const PtBufferView &view,
+spu::Value Constant(SPUContext *ctx, const PtBufferView &view,
                     absl::Span<const int64_t> out_shape) {
   const auto dtype = getEncodeType(view.pt_type);
   if (view.shape == out_shape) {
@@ -32,10 +32,10 @@ spu::Value Constant(HalContext *ctx, const PtBufferView &view,
   }
 }
 
-spu::Value Iota(HalContext *ctx, DataType dtype, int64_t numel) {
+spu::Value Iota(SPUContext *ctx, DataType dtype, int64_t numel) {
   return hal::iota(ctx, dtype, numel);
 }
 
-spu::Value Epsilon(HalContext *ctx) { return hal::epsilon(ctx); }
+spu::Value Epsilon(SPUContext *ctx) { return hal::epsilon(ctx); }
 
 }  // namespace spu::kernel::hlo
diff --git a/libspu/kernel/hlo/const.h b/libspu/kernel/hlo/const.h
index 48800c7a..951b2ca5 100644
--- a/libspu/kernel/hlo/const.h
+++ b/libspu/kernel/hlo/const.h
@@ -22,14 +22,14 @@
 
 namespace spu::kernel::hlo {
 
-spu::Value Constant(HalContext *ctx, const PtBufferView &view,
+spu::Value Constant(SPUContext *ctx, const PtBufferView &view,
                     absl::Span<const int64_t> out_shape);
 
-spu::Value Iota(HalContext *ctx, DataType dtype, int64_t numel);
+spu::Value Iota(SPUContext *ctx, DataType dtype, int64_t numel);
 
 // Returns the SPU epsilon, the positive distance between two fixed point value.
 //
 // The result visibility is public.
-spu::Value Epsilon(HalContext *ctx);
+spu::Value Epsilon(SPUContext *ctx);
 
 }  // namespace spu::kernel::hlo
diff --git a/libspu/kernel/hlo/control_flow.cc b/libspu/kernel/hlo/control_flow.cc
index 4258092a..5dfb8c02 100644
--- a/libspu/kernel/hlo/control_flow.cc
+++ b/libspu/kernel/hlo/control_flow.cc
@@ -28,7 +28,7 @@
 
 namespace spu::kernel::hlo {
 
-std::vector<spu::Value> IfElse(HalContext *ctx, const spu::Value &condition,
+std::vector<spu::Value> IfElse(SPUContext *ctx, const spu::Value &condition,
                                const BranchFcnT &on_true,
                                const BranchFcnT &on_false) {
   if (condition.isSecret()) {
@@ -51,7 +51,7 @@ std::vector<spu::Value> IfElse(HalContext *ctx, const spu::Value &condition,
   }
 }
 
-std::vector<spu::Value> Case(HalContext *ctx, const spu::Value &index,
+std::vector<spu::Value> Case(SPUContext *ctx, const spu::Value &index,
                              absl::Span<const BranchFcnT> branches) {
   SPU_ENFORCE(index.isInt());
   if (index.isPublic()) {
@@ -112,7 +112,7 @@ std::vector<spu::Value> Case(HalContext *ctx, const spu::Value &index,
   }
 }
 
-std::vector<spu::Value> While(HalContext *ctx,
+std::vector<spu::Value> While(SPUContext *ctx,
                               absl::Span<const spu::Value> inputs,
                               const ConditionFcnT &cond, const BodyFcnT &body) {
   bool warned = false;
diff --git a/libspu/kernel/hlo/control_flow.h b/libspu/kernel/hlo/control_flow.h
index 920e8953..c829beab 100644
--- a/libspu/kernel/hlo/control_flow.h
+++ b/libspu/kernel/hlo/control_flow.h
@@ -14,18 +14,18 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hlo {
 
 using BranchFcnT = std::function<std::vector<spu::Value>()>;
 
-std::vector<spu::Value> IfElse(HalContext *ctx, const spu::Value &condition,
+std::vector<spu::Value> IfElse(SPUContext *ctx, const spu::Value &condition,
                                const BranchFcnT &on_true,
                                const BranchFcnT &on_false);
 
-std::vector<spu::Value> Case(HalContext *ctx, const spu::Value &index,
+std::vector<spu::Value> Case(SPUContext *ctx, const spu::Value &index,
                              absl::Span<const BranchFcnT> branches);
 
 /// While evaluation order:
@@ -36,7 +36,7 @@ std::vector<spu::Value> Case(HalContext *ctx, const spu::Value &index,
 using ConditionFcnT = std::function<spu::Value(absl::Span<const spu::Value>)>;
 using BodyFcnT =
     std::function<std::vector<spu::Value>(absl::Span<const spu::Value>)>;
-std::vector<spu::Value> While(HalContext *ctx,
+std::vector<spu::Value> While(SPUContext *ctx,
                               absl::Span<const spu::Value> inputs,
                               const ConditionFcnT &cond, const BodyFcnT &body);
 
diff --git a/libspu/kernel/hlo/convolution.cc b/libspu/kernel/hlo/convolution.cc
index a97383d5..312009bb 100644
--- a/libspu/kernel/hlo/convolution.cc
+++ b/libspu/kernel/hlo/convolution.cc
@@ -14,13 +14,13 @@
 
 #include "libspu/kernel/hlo/convolution.h"
 
-#include "libspu/kernel/context.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hal/constants.h"
 #include "libspu/kernel/hal/polymorphic.h"
 #include "libspu/kernel/hal/shape_ops.h"
 #include "libspu/kernel/hal/type_cast.h"
 #include "libspu/kernel/hlo/utils.h"
-#include "libspu/kernel/value.h"
 
 namespace {
 
@@ -38,7 +38,7 @@ std::vector<int64_t> MakeDimMultipliers(absl::Span<const int64_t> shape) {
 
 namespace spu::kernel::hlo {
 
-spu::Value Convolution(HalContext *ctx, const spu::Value &lhs,
+spu::Value Convolution(SPUContext *ctx, const spu::Value &lhs,
                        const spu::Value &rhs, const ConvolutionConfig &config,
                        absl::Span<const int64_t> result_shape) {
   const size_t num_spatial_dims = config.outputSpatialDimensions.size();
@@ -186,7 +186,7 @@ spu::Value Convolution(HalContext *ctx, const spu::Value &lhs,
   return ret;
 }
 
-spu::Value extractImagePatches(HalContext *ctx, spu::Value &input,
+spu::Value extractImagePatches(SPUContext *ctx, spu::Value &input,
                                int64_t kernel_x, int64_t kernel_y,
                                int64_t stride_x, int64_t stride_y) {
   auto input_batch = input.shape()[0];
@@ -214,7 +214,7 @@ spu::Value extractImagePatches(HalContext *ctx, spu::Value &input,
 }
 
 // This is an optimized conv2D with im2col
-spu::Value Convolution2D(HalContext *ctx, spu::Value input,
+spu::Value Convolution2D(SPUContext *ctx, spu::Value input,
                          const spu::Value &kernel,
                          const ConvolutionConfig &config,
                          absl::Span<const int64_t> result_shape) {
@@ -228,8 +228,8 @@ spu::Value Convolution2D(HalContext *ctx, spu::Value input,
   auto output_x = result_shape[1];
   auto output_y = result_shape[2];
 
-  if (ctx->rt_config().protocol() == ProtocolKind::CHEETAH &&
-      input.isSecret() && kernel.isSecret()) {
+  if (ctx->config().protocol() == ProtocolKind::CHEETAH && input.isSecret() &&
+      kernel.isSecret()) {
     // NOTE(juhou): ad-hoc optimization for the current 2PC conv2d
     // implementation. When input_batch is large or small kernel size, it would
     // be better to compute im2col because the current conv2d implementation
diff --git a/libspu/kernel/hlo/convolution.h b/libspu/kernel/hlo/convolution.h
index 2b92cb61..283939d5 100644
--- a/libspu/kernel/hlo/convolution.h
+++ b/libspu/kernel/hlo/convolution.h
@@ -36,12 +36,12 @@ struct ConvolutionConfig {
 // This is a port of hlo evaluator's HandleConvolutionWithLiterals, which can
 // handle general convolution. See
 // https://github.com/tensorflow/tensorflow/blob/master/tensorflow/compiler/xla/service/hlo_evaluator_typed_visitor.h
-spu::Value Convolution(HalContext *ctx, const spu::Value &lhs,
+spu::Value Convolution(SPUContext *ctx, const spu::Value &lhs,
                        const spu::Value &rhs, const ConvolutionConfig &config,
                        absl::Span<const int64_t> result_shape);
 
 // This is an optimized conv2D with im2col
-spu::Value Convolution2D(HalContext *ctx, spu::Value input,
+spu::Value Convolution2D(SPUContext *ctx, spu::Value input,
                          const spu::Value &kernel,
                          const ConvolutionConfig &config,
                          absl::Span<const int64_t> result_shape);
diff --git a/libspu/kernel/hlo/geometrical.cc b/libspu/kernel/hlo/geometrical.cc
index f33017c2..f87da8e4 100644
--- a/libspu/kernel/hlo/geometrical.cc
+++ b/libspu/kernel/hlo/geometrical.cc
@@ -18,41 +18,41 @@
 
 namespace spu::kernel::hlo {
 
-spu::Value Transpose(HalContext *ctx, const spu::Value &in,
+spu::Value Transpose(SPUContext *ctx, const spu::Value &in,
                      absl::Span<const int64_t> permutation) {
   return hal::transpose(ctx, in, permutation);
 }
 
-spu::Value Broadcast(HalContext *ctx, const spu::Value &in,
+spu::Value Broadcast(SPUContext *ctx, const spu::Value &in,
                      absl::Span<const int64_t> to_shape,
                      absl::Span<const int64_t> in_dims) {
   return hal::broadcast_to(ctx, in, to_shape, in_dims);
 }
 
-spu::Value Reshape(HalContext *ctx, const spu::Value &in,
+spu::Value Reshape(SPUContext *ctx, const spu::Value &in,
                    absl::Span<const int64_t> to_shape) {
   return hal::reshape(ctx, in, to_shape);
 }
 
-spu::Value Concatenate(HalContext *ctx, absl::Span<const spu::Value> operands,
+spu::Value Concatenate(SPUContext *ctx, absl::Span<const spu::Value> operands,
                        int64_t axis) {
   return hal::concatenate(ctx, operands, axis);
 }
 
-spu::Value Slice(HalContext *ctx, const spu::Value &in,
+spu::Value Slice(SPUContext *ctx, const spu::Value &in,
                  absl::Span<const int64_t> start, absl::Span<const int64_t> end,
                  absl::Span<const int64_t> strides) {
   return hal::slice(ctx, in, start, end, strides);
 }
 
-spu::Value Pad(HalContext *ctx, const spu::Value &in,
+spu::Value Pad(SPUContext *ctx, const spu::Value &in,
                const spu::Value &pad_value, absl::Span<const int64_t> edge_low,
                absl::Span<const int64_t> edge_high,
                absl::Span<const int64_t> inner) {
   return hal::pad(ctx, in, pad_value, edge_low, edge_high, inner);
 }
 
-spu::Value Reverse(HalContext *ctx, const spu::Value &in,
+spu::Value Reverse(SPUContext *ctx, const spu::Value &in,
                    absl::Span<const int64_t> dims) {
   return hal::reverse(ctx, in, dims);
 }
diff --git a/libspu/kernel/hlo/geometrical.h b/libspu/kernel/hlo/geometrical.h
index bd5852c0..f954eae6 100644
--- a/libspu/kernel/hlo/geometrical.h
+++ b/libspu/kernel/hlo/geometrical.h
@@ -17,29 +17,29 @@
 #include "libspu/kernel/hlo/utils.h"
 namespace spu::kernel::hlo {
 
-spu::Value Transpose(HalContext *ctx, const spu::Value &in,
+spu::Value Transpose(SPUContext *ctx, const spu::Value &in,
                      absl::Span<const int64_t> permutation);
 
-spu::Value Broadcast(HalContext *ctx, const spu::Value &in,
+spu::Value Broadcast(SPUContext *ctx, const spu::Value &in,
                      absl::Span<const int64_t> to_shape,
                      absl::Span<const int64_t> in_dims);
 
-spu::Value Reshape(HalContext *ctx, const spu::Value &in,
+spu::Value Reshape(SPUContext *ctx, const spu::Value &in,
                    absl::Span<const int64_t> to_shape);
 
-spu::Value Concatenate(HalContext *ctx, absl::Span<const spu::Value> operands,
+spu::Value Concatenate(SPUContext *ctx, absl::Span<const spu::Value> operands,
                        int64_t axis);
 
-spu::Value Slice(HalContext *ctx, const spu::Value &in,
+spu::Value Slice(SPUContext *ctx, const spu::Value &in,
                  absl::Span<const int64_t> start, absl::Span<const int64_t> end,
                  absl::Span<const int64_t> strides);
 
-spu::Value Pad(HalContext *ctx, const spu::Value &in,
+spu::Value Pad(SPUContext *ctx, const spu::Value &in,
                const spu::Value &pad_value, absl::Span<const int64_t> edge_low,
                absl::Span<const int64_t> edge_high,
                absl::Span<const int64_t> inner);
 
-spu::Value Reverse(HalContext *ctx, const spu::Value &in,
+spu::Value Reverse(SPUContext *ctx, const spu::Value &in,
                    absl::Span<const int64_t> dims);
 
 }  // namespace spu::kernel::hlo
diff --git a/libspu/kernel/hlo/indexing.cc b/libspu/kernel/hlo/indexing.cc
index cb5faace..9d7df570 100644
--- a/libspu/kernel/hlo/indexing.cc
+++ b/libspu/kernel/hlo/indexing.cc
@@ -19,6 +19,7 @@
 #include "llvm/ADT/STLExtras.h"
 
 #include "libspu/core/ndarray_ref.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hal/hal.h"
 #include "libspu/kernel/hal/ring.h"
 #include "libspu/kernel/hlo/basic_binary.h"
@@ -28,7 +29,6 @@
 #include "libspu/kernel/hlo/geometrical.h"
 #include "libspu/kernel/hlo/reduce.h"
 #include "libspu/kernel/hlo/utils.h"
-#include "libspu/kernel/value.h"
 
 namespace {
 struct IndexIterationSpace {
@@ -294,7 +294,7 @@ class OutputOffsetIndexToInputIndex {
   std::vector<int64_t> input_index_;
 };
 
-spu::Value reshapedGatherIndices(spu::HalContext *ctx, int64_t index_vector_dim,
+spu::Value reshapedGatherIndices(spu::SPUContext *ctx, int64_t index_vector_dim,
                                  const spu::Value &start_indices) {
   if (start_indices.shape().size() != static_cast<size_t>(index_vector_dim)) {
     return start_indices;
@@ -306,7 +306,7 @@ spu::Value reshapedGatherIndices(spu::HalContext *ctx, int64_t index_vector_dim,
   return spu::kernel::hal::reshape(ctx, start_indices, new_shape);
 }
 
-spu::Value SecretLinearUpdateIndexing(spu::HalContext *ctx,
+spu::Value SecretLinearUpdateIndexing(spu::SPUContext *ctx,
                                       const spu::Value &operand,
                                       const spu::Value &update,
                                       const spu::Value &linear_idx) {
@@ -337,7 +337,7 @@ spu::Value SecretLinearUpdateIndexing(spu::HalContext *ctx,
 }
 
 std::vector<spu::Value> ClampAndFlattenIndex(
-    spu::HalContext *ctx, absl::Span<const spu::Value> start_indices,
+    spu::SPUContext *ctx, absl::Span<const spu::Value> start_indices,
     absl::Span<const int64_t> iterate_shape,
     absl::Span<const int64_t> limit_shape) {
   // Transform start_indices
@@ -420,7 +420,7 @@ std::vector<spu::Value> ClampAndFlattenIndex(
 
 namespace spu::kernel::hlo {
 
-spu::Value Gather(HalContext *ctx, const spu::Value &operand,
+spu::Value Gather(SPUContext *ctx, const spu::Value &operand,
                   const spu::Value &start_indices, const GatherConfig &config,
                   absl::Span<const int64_t> result_shape) {
   // If input is empty, short circuit
@@ -517,7 +517,7 @@ spu::Value Gather(HalContext *ctx, const spu::Value &operand,
   return result;
 }
 
-spu::Value DynamicUpdateSlice(HalContext *ctx, const spu::Value &operand,
+spu::Value DynamicUpdateSlice(SPUContext *ctx, const spu::Value &operand,
                               const spu::Value &update,
                               absl::Span<const spu::Value> start_indices) {
   // Basic idea here, get a ref slice and update the whole slice..
@@ -569,14 +569,14 @@ spu::Value DynamicUpdateSlice(HalContext *ctx, const spu::Value &operand,
   }
 }
 
-spu::Value UpdateSlice(HalContext *ctx, const spu::Value &in,
+spu::Value UpdateSlice(SPUContext *ctx, const spu::Value &in,
                        const spu::Value &update,
                        absl::Span<const int64_t> start_indices) {
   return hal::update_slice(ctx, in, update, start_indices);
 }
 
 spu::Value SecretDynamicSlice(
-    HalContext *ctx, const spu::Value &operand,
+    SPUContext *ctx, const spu::Value &operand,
     absl::Span<const int64_t> slice_size,
     absl::Span<const spu::Value> start_indices,
     std::vector<std::optional<spu::Value>>::iterator index_cache_mask) {
@@ -677,7 +677,7 @@ spu::Value SecretDynamicSlice(
   return hal::concatenate(ctx, results, 0);
 }
 
-spu::Value DynamicSlice(HalContext *ctx, const spu::Value &operand,
+spu::Value DynamicSlice(SPUContext *ctx, const spu::Value &operand,
                         absl::Span<const int64_t> slice_size,
                         absl::Span<const spu::Value> start_indices) {
   SPU_ENFORCE_EQ(slice_size.size(), start_indices.size());
@@ -716,7 +716,7 @@ spu::Value DynamicSlice(HalContext *ctx, const spu::Value &operand,
   }
 }
 
-spu::Value FilterByMask(HalContext *ctx, const spu::Value &operand,
+spu::Value FilterByMask(SPUContext *ctx, const spu::Value &operand,
                         absl::Span<const uint8_t> mask) {
   // Sanity
   SPU_ENFORCE(operand.shape().size() == 1, "Operand must be a vector");
@@ -743,12 +743,12 @@ spu::Value FilterByMask(HalContext *ctx, const spu::Value &operand,
   return Value(operand.data().linear_gather(indices), operand.dtype());
 }
 
-spu::Value LinearGather(HalContext *ctx, const spu::Value &in,
+spu::Value LinearGather(SPUContext *ctx, const spu::Value &in,
                         absl::Span<const int64_t> indices) {
   return Value(in.data().linear_gather(indices), in.dtype());
 }
 
-void LinearScatterInPlace(HalContext *ctx, spu::Value &in,
+void LinearScatterInPlace(SPUContext *ctx, spu::Value &in,
                           const spu::Value &update,
                           absl::Span<const int64_t> indices) {
   if (in.data().eltype() != update.data().eltype()) {
diff --git a/libspu/kernel/hlo/indexing.h b/libspu/kernel/hlo/indexing.h
index d12852e4..ca3cb0d6 100644
--- a/libspu/kernel/hlo/indexing.h
+++ b/libspu/kernel/hlo/indexing.h
@@ -14,7 +14,7 @@
 
 #pragma once
 
-#include "libspu/kernel/context.h"
+#include "libspu/core/context.h"
 #include "libspu/kernel/hlo/utils.h"
 
 namespace spu::kernel::hlo {
@@ -29,31 +29,31 @@ struct GatherConfig {
 
 // This is ported from
 // https://github.com/tensorflow/tensorflow/blob/bf4c6ad46dac1f7f69911e2bfc48e141a39b40af/tensorflow/compiler/xla/service/hlo_evaluator.cc#L1774
-spu::Value Gather(HalContext *ctx, const spu::Value &operand,
+spu::Value Gather(SPUContext *ctx, const spu::Value &operand,
                   const spu::Value &start_indices, const GatherConfig &config,
                   absl::Span<const int64_t> result_shape);
 
-spu::Value DynamicUpdateSlice(HalContext *ctx, const spu::Value &operand,
+spu::Value DynamicUpdateSlice(SPUContext *ctx, const spu::Value &operand,
                               const spu::Value &update,
                               absl::Span<const spu::Value> start_indices);
 
-spu::Value DynamicSlice(HalContext *ctx, const spu::Value &operand,
+spu::Value DynamicSlice(SPUContext *ctx, const spu::Value &operand,
                         absl::Span<const int64_t> slice_size,
                         absl::Span<const spu::Value> start_indices);
 
 /// ------------------- non-XLA APIs ------------------------------------
 // @brief Update slice
-spu::Value UpdateSlice(HalContext *ctx, const spu::Value &in,
+spu::Value UpdateSlice(SPUContext *ctx, const spu::Value &in,
                        const spu::Value &update,
                        absl::Span<const int64_t> start_indices);
 
-spu::Value FilterByMask(HalContext *ctx, const spu::Value &operand,
+spu::Value FilterByMask(SPUContext *ctx, const spu::Value &operand,
                         absl::Span<const uint8_t> mask);
 
-spu::Value LinearGather(HalContext *ctx, const spu::Value &in,
+spu::Value LinearGather(SPUContext *ctx, const spu::Value &in,
                         absl::Span<const int64_t> indices);
 
-void LinearScatterInPlace(HalContext *ctx, spu::Value &in,
+void LinearScatterInPlace(SPUContext *ctx, spu::Value &in,
                           const spu::Value &update,
                           absl::Span<const int64_t> indices);
 
diff --git a/libspu/kernel/hlo/rand.cc b/libspu/kernel/hlo/rand.cc
index 21a591be..795e2348 100644
--- a/libspu/kernel/hlo/rand.cc
+++ b/libspu/kernel/hlo/rand.cc
@@ -18,7 +18,7 @@
 
 namespace spu::kernel::hlo {
 
-spu::Value Uniform_rand(HalContext *ctx, const spu::Value &lo,
+spu::Value Uniform_rand(SPUContext *ctx, const spu::Value &lo,
                         const spu::Value &hi,
                         absl::Span<const int64_t> to_shape) {
   return hal::rng_uniform(ctx, lo, hi, to_shape);
diff --git a/libspu/kernel/hlo/rand.h b/libspu/kernel/hlo/rand.h
index 82c46795..7e77b0eb 100644
--- a/libspu/kernel/hlo/rand.h
+++ b/libspu/kernel/hlo/rand.h
@@ -18,7 +18,7 @@
 
 namespace spu::kernel::hlo {
 
-spu::Value Uniform_rand(HalContext *ctx, const spu::Value &lo,
+spu::Value Uniform_rand(SPUContext *ctx, const spu::Value &lo,
                         const spu::Value &hi,
                         absl::Span<const int64_t> to_shape);
 
diff --git a/libspu/kernel/hlo/reduce.cc b/libspu/kernel/hlo/reduce.cc
index 305d4604..a4d54bdb 100644
--- a/libspu/kernel/hlo/reduce.cc
+++ b/libspu/kernel/hlo/reduce.cc
@@ -34,7 +34,7 @@
 
 namespace spu::kernel::hlo {
 
-std::vector<spu::Value> TreeReduce(HalContext *ctx,
+std::vector<spu::Value> TreeReduce(SPUContext *ctx,
                                    absl::Span<const spu::Value> inputs,
                                    int64_t axis,
                                    const BatchedValueBinaryFn &reducer) {
@@ -110,7 +110,7 @@ std::vector<spu::Value> TreeReduce(HalContext *ctx,
 }
 
 spu::Value ExpandStridedWindow(
-    HalContext *ctx, const spu::Value &base,
+    SPUContext *ctx, const spu::Value &base,
     absl::Span<const int64_t> window_shape,
     absl::Span<const int64_t> window_strides,
     absl::Span<const std::pair<int64_t, int64_t>> padding) {
@@ -177,7 +177,7 @@ spu::Value ExpandStridedWindow(
   return expanded;
 }
 
-spu::Value ConvertToTiledLayout(HalContext *ctx, const spu::Value &in,
+spu::Value ConvertToTiledLayout(SPUContext *ctx, const spu::Value &in,
                                 absl::Span<const int64_t> block_shape) {
   // Note(jint): use pad+reshape+transpose to convert from column layout to
   // tiled layout.
@@ -201,7 +201,7 @@ spu::Value ConvertToTiledLayout(HalContext *ctx, const spu::Value &in,
 }
 
 std::vector<spu::Value> ReduceWindowWithoutDilation(
-    HalContext *ctx, absl::Span<const spu::Value> inputs,
+    SPUContext *ctx, absl::Span<const spu::Value> inputs,
     absl::Span<const spu::Value> init_values,
     absl::Span<const int64_t> window_shape,
     absl::Span<const int64_t> window_strides,
@@ -284,7 +284,7 @@ std::vector<spu::Value> ReduceWindowWithoutDilation(
 }
 
 std::vector<spu::Value> ReduceWindowImpl(
-    HalContext *ctx, absl::Span<const spu::Value> inputs,
+    SPUContext *ctx, absl::Span<const spu::Value> inputs,
     absl::Span<const spu::Value> init_values,
     absl::Span<const int64_t> ret_shape, const ReduceWindowConfig &config,
     bool last_operand_is_window_mask, bool ignore_init_value,
@@ -357,7 +357,7 @@ std::vector<spu::Value> ReduceWindowImpl(
   return rets;
 }
 
-std::vector<spu::Value> ReduceWindow(HalContext *ctx,
+std::vector<spu::Value> ReduceWindow(SPUContext *ctx,
                                      absl::Span<const spu::Value> inputs,
                                      absl::Span<const spu::Value> init_values,
                                      absl::Span<const int64_t> ret_shape,
@@ -368,7 +368,7 @@ std::vector<spu::Value> ReduceWindow(HalContext *ctx,
                           ignore_init_values, reducer);
 }
 
-std::vector<spu::Value> Reduce(HalContext *ctx,
+std::vector<spu::Value> Reduce(SPUContext *ctx,
                                absl::Span<const spu::Value> inputs,
                                absl::Span<const spu::Value> init_values,
                                absl::Span<const int64_t> dims_to_reduce,
@@ -465,7 +465,7 @@ std::vector<spu::Value> Reduce(HalContext *ctx,
 // And without dilation and padding, this can be achieved through just slicing
 // FIXME: This is a super special case...consider generalize it a little bit
 std::pair<spu::Value, spu::Value> ArgMax1x2x2x1NoPaddingWithoutDilation(
-    HalContext *ctx, const spu::Value &input,
+    SPUContext *ctx, const spu::Value &input,
     absl::Span<const int64_t> window_strides) {
   auto input_shape = input.shape();
 
@@ -564,7 +564,7 @@ std::pair<spu::Value, spu::Value> ArgMax1x2x2x1NoPaddingWithoutDilation(
   return {max_ret, max_indices};
 }
 
-std::pair<spu::Value, spu::Value> ArgMax(HalContext *ctx,
+std::pair<spu::Value, spu::Value> ArgMax(SPUContext *ctx,
                                          const spu::Value &input,
                                          absl::Span<const int64_t> ret_shape,
                                          const ReduceWindowConfig &config) {
diff --git a/libspu/kernel/hlo/reduce.h b/libspu/kernel/hlo/reduce.h
index 8b79722e..dc2edf67 100644
--- a/libspu/kernel/hlo/reduce.h
+++ b/libspu/kernel/hlo/reduce.h
@@ -19,8 +19,8 @@
 
 #include "absl/types/span.h"
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hlo {
 
@@ -28,15 +28,15 @@ using BatchedValueBinaryFn = std::function<std::vector<spu::Value>(
     absl::Span<spu::Value const> lhs, absl::Span<spu::Value const> rhs)>;
 
 spu::Value ExpandStridedWindow(
-    HalContext *ctx, const spu::Value &base,
+    SPUContext *ctx, const spu::Value &base,
     absl::Span<const int64_t> window_shape,
     absl::Span<const int64_t> window_strides,
     absl::Span<const std::pair<int64_t, int64_t>> padding);
 
-spu::Value ConvertToTiledLayout(HalContext *ctx, const spu::Value &in,
+spu::Value ConvertToTiledLayout(SPUContext *ctx, const spu::Value &in,
                                 absl::Span<const int64_t> block_shape);
 
-std::vector<spu::Value> TreeReduce(HalContext *ctx,
+std::vector<spu::Value> TreeReduce(SPUContext *ctx,
                                    absl::Span<const spu::Value> inputs,
                                    int64_t axis,
                                    const BatchedValueBinaryFn &reducer);
@@ -49,7 +49,7 @@ struct ReduceWindowConfig {
   absl::Span<const int64_t> base_dilations;
 };
 
-std::vector<spu::Value> ReduceWindow(HalContext *ctx,
+std::vector<spu::Value> ReduceWindow(SPUContext *ctx,
                                      absl::Span<const spu::Value> inputs,
                                      absl::Span<const spu::Value> init_values,
                                      absl::Span<const int64_t> ret_shape,
@@ -57,14 +57,14 @@ std::vector<spu::Value> ReduceWindow(HalContext *ctx,
                                      const BatchedValueBinaryFn &reducer,
                                      bool ignore_init_values = false);
 
-std::vector<spu::Value> Reduce(HalContext *ctx,
+std::vector<spu::Value> Reduce(SPUContext *ctx,
                                absl::Span<const spu::Value> inputs,
                                absl::Span<const spu::Value> init_values,
                                absl::Span<const int64_t> dims_to_reduce,
                                const BatchedValueBinaryFn &reducer,
                                bool ignore_init_values = false);
 
-std::pair<spu::Value, spu::Value> ArgMax(HalContext *ctx,
+std::pair<spu::Value, spu::Value> ArgMax(SPUContext *ctx,
                                          const spu::Value &input,
                                          absl::Span<const int64_t> ret_shape,
                                          const ReduceWindowConfig &config);
diff --git a/libspu/kernel/hlo/select_and_scatter.cc b/libspu/kernel/hlo/select_and_scatter.cc
index 4db69690..e3d1042e 100644
--- a/libspu/kernel/hlo/select_and_scatter.cc
+++ b/libspu/kernel/hlo/select_and_scatter.cc
@@ -22,8 +22,9 @@
 
 #include "yacl/utils/parallel.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/shape_util.h"
-#include "libspu/kernel/context.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hal/constants.h"
 #include "libspu/kernel/hal/debug.h"
 #include "libspu/kernel/hal/polymorphic.h"  // for select
@@ -33,12 +34,11 @@
 #include "libspu/kernel/hlo/const.h"
 #include "libspu/kernel/hlo/reduce.h"
 #include "libspu/kernel/hlo/utils.h"
-#include "libspu/kernel/value.h"
 
 namespace spu::kernel::hlo {
 
 spu::Value MaxPoolScatter1x2x2x1NoPaddingNoDialation(
-    HalContext *ctx, const spu::Value &scatter_indices,
+    SPUContext *ctx, const spu::Value &scatter_indices,
     const spu::Value &source, absl::Span<const int64_t> window_strides) {
   std::vector<spu::Value> slices(4);
   for (int64_t idx = 0; idx < 4; ++idx) {
@@ -88,7 +88,7 @@ spu::Value MaxPoolScatter1x2x2x1NoPaddingNoDialation(
 };
 
 spu::Value MaxPoolScatter(
-    HalContext *ctx, const spu::Value &scatter_indices,
+    SPUContext *ctx, const spu::Value &scatter_indices,
     const spu::Value &source, absl::Span<const int64_t> window_shape,
     absl::Span<const int64_t> base_shape,
     absl::Span<const int64_t> window_strides,
@@ -180,7 +180,7 @@ spu::Value MaxPoolScatter(
 }
 
 spu::Value SelectAndScatterExpanded(
-    HalContext *ctx, const spu::Value &base, const spu::Value &source,
+    SPUContext *ctx, const spu::Value &base, const spu::Value &source,
     const spu::Value &init_val, absl::Span<const int64_t> window_shape,
     absl::Span<const int64_t> window_strides,
     absl::Span<const std::pair<int64_t, int64_t>> window_padding,
@@ -289,7 +289,7 @@ spu::Value SelectAndScatterExpanded(
 }
 
 spu::Value SelectAndScatterNaive(
-    HalContext *ctx, const spu::Value &operand, const spu::Value &source,
+    SPUContext *ctx, const spu::Value &operand, const spu::Value &source,
     const spu::Value &init_val, absl::Span<const int64_t> window_shape,
     absl::Span<const int64_t> window_strides,
     absl::Span<const std::pair<int64_t, int64_t>> window_padding,
diff --git a/libspu/kernel/hlo/select_and_scatter.h b/libspu/kernel/hlo/select_and_scatter.h
index 959ec0b2..816bdb8a 100644
--- a/libspu/kernel/hlo/select_and_scatter.h
+++ b/libspu/kernel/hlo/select_and_scatter.h
@@ -19,8 +19,8 @@
 
 #include "absl/types/span.h"
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hlo {
 
@@ -28,14 +28,14 @@ using ValueBinaryFn =
     std::function<spu::Value(spu::Value const &lhs, spu::Value const &rhs)>;
 
 spu::Value SelectAndScatterExpanded(
-    HalContext *ctx, const spu::Value &base, const spu::Value &source,
+    SPUContext *ctx, const spu::Value &base, const spu::Value &source,
     const spu::Value &init_val, absl::Span<const int64_t> window_shape,
     absl::Span<const int64_t> window_strides,
     absl::Span<const std::pair<int64_t, int64_t>> window_padding,
     const ValueBinaryFn &select_fn, const ValueBinaryFn &scatter_fn);
 
 spu::Value SelectAndScatterNaive(
-    HalContext *ctx, const spu::Value &operand, const spu::Value &source,
+    SPUContext *ctx, const spu::Value &operand, const spu::Value &source,
     const spu::Value &init_val, absl::Span<const int64_t> window_shape,
     absl::Span<const int64_t> window_strides,
     absl::Span<const std::pair<int64_t, int64_t>> window_padding,
@@ -54,7 +54,7 @@ spu::Value SelectAndScatterNaive(
  * @return spu::Value
  */
 spu::Value MaxPoolScatter(
-    HalContext *ctx, const spu::Value &scatter_indices,
+    SPUContext *ctx, const spu::Value &scatter_indices,
     const spu::Value &source, absl::Span<const int64_t> window_shape,
     absl::Span<const int64_t> base_shape,
     absl::Span<const int64_t> window_strides,
diff --git a/libspu/kernel/hlo/shift.cc b/libspu/kernel/hlo/shift.cc
index b84bc731..bc3c76e5 100644
--- a/libspu/kernel/hlo/shift.cc
+++ b/libspu/kernel/hlo/shift.cc
@@ -25,7 +25,7 @@
 namespace spu::kernel::hlo {
 
 template <typename Fn>
-spu::Value shift_impl_p(HalContext *ctx, const spu::Value &lhs,
+spu::Value shift_impl_p(SPUContext *ctx, const spu::Value &lhs,
                         const spu::Value &rhs, const Fn &f) {
   auto shift_bits = hal::dump_public_as<int8_t>(ctx, rhs);
   if (std::all_of(rhs.strides().begin(), rhs.strides().end(),
@@ -55,7 +55,7 @@ spu::Value shift_impl_p(HalContext *ctx, const spu::Value &lhs,
 }
 
 template <typename Fn>
-spu::Value shift_impl_s(HalContext *ctx, const spu::Value &lhs,
+spu::Value shift_impl_s(SPUContext *ctx, const spu::Value &lhs,
                         const spu::Value &rhs, const Fn &f) {
   spu::Value ret =
       hal::constant(ctx, static_cast<uint8_t>(0), lhs.dtype(), lhs.shape());
@@ -72,7 +72,7 @@ spu::Value shift_impl_s(HalContext *ctx, const spu::Value &lhs,
 }
 
 template <typename Fn>
-spu::Value shift_impl(HalContext *ctx, const spu::Value &lhs,
+spu::Value shift_impl(SPUContext *ctx, const spu::Value &lhs,
                       const spu::Value &rhs, const Fn &f) {
   SPU_ENFORCE(rhs.shape() == lhs.shape());
 
@@ -83,17 +83,17 @@ spu::Value shift_impl(HalContext *ctx, const spu::Value &lhs,
   return shift_impl_s(ctx, lhs, rhs, f);
 }
 
-spu::Value Lshift(HalContext *ctx, const spu::Value &operand,
+spu::Value Lshift(SPUContext *ctx, const spu::Value &operand,
                   const spu::Value &bits_to_shift) {
   return shift_impl(ctx, operand, bits_to_shift, hal::left_shift);
 }
 
-spu::Value ARshift(HalContext *ctx, const spu::Value &operand,
+spu::Value ARshift(SPUContext *ctx, const spu::Value &operand,
                    const spu::Value &bits_to_shift) {
   return shift_impl(ctx, operand, bits_to_shift, hal::right_shift_arithmetic);
 }
 
-spu::Value Rshift(HalContext *ctx, const spu::Value &operand,
+spu::Value Rshift(SPUContext *ctx, const spu::Value &operand,
                   const spu::Value &bits_to_shift) {
   return shift_impl(ctx, operand, bits_to_shift, hal::right_shift_logical);
 }
diff --git a/libspu/kernel/hlo/shift.h b/libspu/kernel/hlo/shift.h
index 5d23562c..ea2e6bcb 100644
--- a/libspu/kernel/hlo/shift.h
+++ b/libspu/kernel/hlo/shift.h
@@ -18,13 +18,13 @@
 
 namespace spu::kernel::hlo {
 
-spu::Value Lshift(HalContext *ctx, const spu::Value &operand,
+spu::Value Lshift(SPUContext *ctx, const spu::Value &operand,
                   const spu::Value &bits_to_shift);
 
-spu::Value Rshift(HalContext *ctx, const spu::Value &operand,
+spu::Value Rshift(SPUContext *ctx, const spu::Value &operand,
                   const spu::Value &bits_to_shift);
 
-spu::Value ARshift(HalContext *ctx, const spu::Value &operand,
+spu::Value ARshift(SPUContext *ctx, const spu::Value &operand,
                    const spu::Value &bits_to_shift);
 
 }  // namespace spu::kernel::hlo
diff --git a/libspu/kernel/hlo/shuffle.cc b/libspu/kernel/hlo/shuffle.cc
index f24f78ba..9555e1e9 100644
--- a/libspu/kernel/hlo/shuffle.cc
+++ b/libspu/kernel/hlo/shuffle.cc
@@ -22,7 +22,7 @@
 
 namespace spu::kernel::hlo {
 
-std::vector<spu::Value> Shuffle(HalContext* ctx,
+std::vector<spu::Value> Shuffle(SPUContext* ctx,
                                 absl::Span<const spu::Value> inputs,
                                 int64_t axis) {
   SPU_ENFORCE_GT(inputs.size(), 0U);
diff --git a/libspu/kernel/hlo/shuffle.h b/libspu/kernel/hlo/shuffle.h
index a9be67d6..71a8aa47 100644
--- a/libspu/kernel/hlo/shuffle.h
+++ b/libspu/kernel/hlo/shuffle.h
@@ -16,13 +16,13 @@
 
 #include "absl/types/span.h"
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel::hlo {
 
 // secret shuffle which means the order is kept secret between parties.
-std::vector<spu::Value> Shuffle(HalContext* ctx,
+std::vector<spu::Value> Shuffle(SPUContext* ctx,
                                 absl::Span<const spu::Value> inputs,
                                 int64_t axis = 0);
 
diff --git a/libspu/kernel/hlo/sort.cc b/libspu/kernel/hlo/sort.cc
index 17b3f75f..4774671a 100644
--- a/libspu/kernel/hlo/sort.cc
+++ b/libspu/kernel/hlo/sort.cc
@@ -26,7 +26,7 @@
 namespace spu::kernel::hlo {
 namespace {
 
-Value permute1D(HalContext *ctx, const Value &x,
+Value permute1D(SPUContext *ctx, const Value &x,
                 absl::Span<const int64_t> indices) {
   SPU_ENFORCE(x.shape().size() == 1);
   return Value(x.data().linear_gather(indices), x.dtype());
@@ -42,7 +42,7 @@ void SliceCopy(spu::Value &dst, const spu::Value &src,
 }
 
 // Sort will be inplace, so always make a copy here.
-std::vector<spu::Value> GetValuesToSort(HalContext *ctx,
+std::vector<spu::Value> GetValuesToSort(SPUContext *ctx,
                                         absl::Span<const spu::Value> inputs,
                                         const std::vector<int64_t> &indices,
                                         int64_t sort_dim,
@@ -71,7 +71,7 @@ std::vector<spu::Value> GetValuesToSort(HalContext *ctx,
 using SequenceT =
     std::vector<std::pair<std::vector<int64_t>, std::vector<int64_t>>>;
 
-void CmpSwap(HalContext *ctx, const CompFn &comparator_body,
+void CmpSwap(SPUContext *ctx, const CompFn &comparator_body,
              std::vector<spu::Value> &values_to_sort,
              absl::Span<const int64_t> lhs_indices,
              absl::Span<const int64_t> rhs_indices) {
@@ -162,7 +162,7 @@ void BuildCmpSwapSequence(SequenceT &seq, int64_t numel) {
   SortSequence(seq, 0, numel, true, depth);
 }
 
-void BitonicSort(HalContext *ctx, const CompFn &comparator_body,
+void BitonicSort(SPUContext *ctx, const CompFn &comparator_body,
                  std::vector<spu::Value> &values_to_sort) {
   // Build a sorting network...
   SequenceT sequence;
@@ -178,7 +178,7 @@ void BitonicSort(HalContext *ctx, const CompFn &comparator_body,
 
 }  // namespace
 
-std::vector<spu::Value> Sort(HalContext *ctx,
+std::vector<spu::Value> Sort(SPUContext *ctx,
                              absl::Span<const spu::Value> inputs,
                              int64_t sort_dim, bool is_stable,
                              const CompFn &comparator_body,
diff --git a/libspu/kernel/hlo/sort.h b/libspu/kernel/hlo/sort.h
index b9034031..06dab853 100644
--- a/libspu/kernel/hlo/sort.h
+++ b/libspu/kernel/hlo/sort.h
@@ -19,16 +19,16 @@
 
 #include "absl/types/span.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/ndarray_ref.h"
-#include "libspu/kernel/context.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hal/shape_ops.h"
-#include "libspu/kernel/value.h"
 
 namespace spu::kernel::hlo {
 
 using CompFn = std::function<spu::Value(absl::Span<const spu::Value>)>;
 
-std::vector<spu::Value> Sort(HalContext* ctx,
+std::vector<spu::Value> Sort(SPUContext* ctx,
                              absl::Span<const spu::Value> inputs,
                              int64_t sort_dim, bool is_stable,
                              const CompFn& comparator_body,
diff --git a/libspu/kernel/hlo/utils.cc b/libspu/kernel/hlo/utils.cc
index fee024f4..b25508d4 100644
--- a/libspu/kernel/hlo/utils.cc
+++ b/libspu/kernel/hlo/utils.cc
@@ -18,7 +18,7 @@
 
 namespace spu::kernel {
 
-bool getBooleanValue(HalContext *ctx, const spu::Value &value) {
+bool getBooleanValue(SPUContext *ctx, const spu::Value &value) {
   SPU_ENFORCE(value.numel() == 1, "Condition value must be a scalar tensor.");
   SPU_ENFORCE(value.dtype() == DT_I1, "Expect bool, got {}", value.dtype());
   SPU_ENFORCE(value.isPublic(), "Expect public value");
@@ -27,7 +27,7 @@ bool getBooleanValue(HalContext *ctx, const spu::Value &value) {
   return public_val.front();
 }
 
-int32_t getI32Value(HalContext *ctx, const spu::Value &value) {
+int32_t getI32Value(SPUContext *ctx, const spu::Value &value) {
   SPU_ENFORCE(value.numel() == 1, "Index value must be a scalar tensor.");
   SPU_ENFORCE(value.dtype() == DT_I32, "Expect bool, got {}", value.dtype());
   SPU_ENFORCE(value.isPublic(), "Expect public value");
@@ -36,7 +36,7 @@ int32_t getI32Value(HalContext *ctx, const spu::Value &value) {
   return public_val.front();
 }
 
-xt::xarray<int64_t> getIndices(HalContext *ctx, const spu::Value &value) {
+xt::xarray<int64_t> getIndices(SPUContext *ctx, const spu::Value &value) {
   SPU_ENFORCE(value.isInt(), "indices value must be integers.");
   SPU_ENFORCE(value.isPublic(), "indices value must be public.");
   return kernel::hal::dump_public_as<int64_t>(ctx, value);
diff --git a/libspu/kernel/hlo/utils.h b/libspu/kernel/hlo/utils.h
index bc5d97ea..12b61352 100644
--- a/libspu/kernel/hlo/utils.h
+++ b/libspu/kernel/hlo/utils.h
@@ -19,16 +19,16 @@
 
 #include "xtensor/xarray.hpp"
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::kernel {
 
-bool getBooleanValue(HalContext *ctx, const spu::Value &value);
+bool getBooleanValue(SPUContext *ctx, const spu::Value &value);
 
-int32_t getI32Value(HalContext *ctx, const spu::Value &value);
+int32_t getI32Value(SPUContext *ctx, const spu::Value &value);
 
-xt::xarray<int64_t> getIndices(HalContext *ctx, const spu::Value &value);
+xt::xarray<int64_t> getIndices(SPUContext *ctx, const spu::Value &value);
 
 template <typename FnTy>
 void forEachIndex(absl::Span<const int64_t> shape,
diff --git a/libspu/mpc/BUILD.bazel b/libspu/mpc/BUILD.bazel
index cf8478b9..c1287986 100644
--- a/libspu/mpc/BUILD.bazel
+++ b/libspu/mpc/BUILD.bazel
@@ -57,19 +57,45 @@ spu_cc_library(
     srcs = ["kernel.cc"],
     hdrs = ["kernel.h"],
     deps = [
-        ":object",
+        "//libspu/core:context",
     ],
 )
 
 spu_cc_library(
-    name = "object",
-    srcs = ["object.cc"],
-    hdrs = ["object.h"],
+    name = "ab_api",
+    srcs = ["ab_api.cc"],
+    hdrs = ["ab_api.h"],
     deps = [
-        "//libspu/core:array_ref",
-        "//libspu/core:type",
-        "//libspu/mpc/utils:cexpr",
+        "//libspu/core:context",
+        "//libspu/core:vectorize",
+        "//libspu/mpc/utils:tiling_util",  # TODO: drop bad reference
+    ],
+)
+
+spu_cc_library(
+    name = "api_test_params",
+    hdrs = ["api_test_params.h"],
+    deps = [
+        "//libspu/core:context",
+    ],
+)
+
+spu_cc_library(
+    name = "ab_api_test",
+    testonly = 1,
+    srcs = ["ab_api_test.cc"],
+    hdrs = ["ab_api_test.h"],
+    deps = [
+        ":ab_api",
+        ":api",
+        ":api_test_params",
+        "//libspu/core:context",
+        "//libspu/mpc:kernel",
+        "//libspu/mpc/common:communicator",
+        "//libspu/mpc/utils:simulate",
+        "@com_google_googletest//:gtest",
     ],
+    alwayslink = True,
 )
 
 spu_cc_library(
@@ -77,8 +103,9 @@ spu_cc_library(
     srcs = ["api.cc"],
     hdrs = ["api.h"],
     deps = [
-        ":object",
+        ":ab_api",
         "//libspu/core:array_ref",
+        "//libspu/core:context",
     ],
 )
 
@@ -89,8 +116,8 @@ spu_cc_library(
     hdrs = ["api_test.h"],
     deps = [
         ":api",
-        ":object",
-        "//libspu/mpc/common:api_test_params",
+        "//libspu/core:context",
+        "//libspu/mpc:api_test_params",
         "//libspu/mpc/common:communicator",
         "//libspu/mpc/utils:simulate",
         "@com_google_googletest//:gtest",
diff --git a/libspu/mpc/ab_api.cc b/libspu/mpc/ab_api.cc
new file mode 100644
index 00000000..9a400e19
--- /dev/null
+++ b/libspu/mpc/ab_api.cc
@@ -0,0 +1,430 @@
+// Copyright 2021 Ant Group Co., Ltd.
+//
+// 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 "libspu/mpc/ab_api.h"
+
+#include "libspu/core/bit_utils.h"
+#include "libspu/core/trace.h"
+#include "libspu/mpc/utils/tiling_util.h"
+
+namespace spu::mpc {
+
+#define FORCE_DISPATCH(CTX, ...)                      \
+  {                                                   \
+    SPU_TRACE_MPC_LEAF(CTX, __VA_ARGS__);             \
+    return dynDispatch((CTX), __func__, __VA_ARGS__); \
+  }
+
+#define TRY_NAMED_DISPATCH(CTX, FNAME, ...)        \
+  if ((CTX)->hasKernel(__func__)) {                \
+    SPU_TRACE_MPC_LEAF(CTX, __VA_ARGS__);          \
+    return dynDispatch((CTX), FNAME, __VA_ARGS__); \
+  }
+
+#define TRY_DISPATCH(CTX, ...) TRY_NAMED_DISPATCH(CTX, __func__, __VA_ARGS__)
+
+template <typename... Args>
+Value tiledDynDispatch(const std::string& fn_name, SPUContext* ctx,
+                       Args&&... args) {
+  auto impl = [fn_name](SPUContext* sh_ctx, Args&&... sh_args) {
+    return dynDispatch(sh_ctx, fn_name, std::forward<Args>(sh_args)...);
+  };
+
+  return tiled(impl, ctx, std::forward<Args>(args)...);
+}
+
+#define TILED_DISPATCH(CTX, ...)                           \
+  {                                                        \
+    SPU_TRACE_MPC_LEAF(ctx, __VA_ARGS__);                  \
+    return tiledDynDispatch(__func__, (CTX), __VA_ARGS__); \
+  }
+
+// TODO: now we handcode mark some of the functions as tiled dispatch according
+// to experience.
+// Note: now tracing is done just before dynamic dispatch, so the time record is
+// OK no matter if tiling is enabled.
+
+Value a2p(SPUContext* ctx, const Value& x) { FORCE_DISPATCH(ctx, x); }
+
+Value p2a(SPUContext* ctx, const Value& x) { FORCE_DISPATCH(ctx, x); }
+
+Value msb_a2b(SPUContext* ctx, const Value& x) { TILED_DISPATCH(ctx, x); }
+
+Value rand_a(SPUContext* ctx, const Shape& shape) {
+  FORCE_DISPATCH(ctx, shape);
+}
+
+Value rand_b(SPUContext* ctx, const Shape& shape) {
+  FORCE_DISPATCH(ctx, shape);
+}
+
+Value not_a(SPUContext* ctx, const Value& x) { FORCE_DISPATCH(ctx, x); }
+
+Value add_ap(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value equal_aa(SPUContext* ctx, const Value& x, const Value& y) {
+  TILED_DISPATCH(ctx, x, y);
+}
+
+Value equal_ap(SPUContext* ctx, const Value& x, const Value& y) {
+  TILED_DISPATCH(ctx, x, y);
+}
+
+Value add_aa(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value mul_ap(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value mul_aa(SPUContext* ctx, const Value& x, const Value& y) {
+  TILED_DISPATCH(ctx, x, y);
+}
+
+Value mul_a1b(SPUContext* ctx, const Value& x, const Value& y) {
+  TILED_DISPATCH(ctx, x, y);
+}
+
+Value lshift_a(SPUContext* ctx, const Value& x, size_t nbits) {
+  FORCE_DISPATCH(ctx, x, nbits);
+}
+
+Value trunc_a(SPUContext* ctx, const Value& x, size_t nbits) {
+  TILED_DISPATCH(ctx, x, nbits);
+}
+
+Value mmul_ap(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k) {
+  FORCE_DISPATCH(ctx, x, y, m, n, k);
+}
+
+Value mmul_aa(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k) {
+  FORCE_DISPATCH(ctx, x, y, m, n, k);
+}
+
+Type common_type_b(SPUContext* ctx, const Type& a, const Type& b) {
+  SPU_TRACE_MPC_LEAF(ctx, a, b);
+  return dynDispatch<Type>(ctx, __func__, a, b);
+}
+
+Value cast_type_b(SPUContext* ctx, const Value& a, const Type& to_type) {
+  FORCE_DISPATCH(ctx, a, to_type);
+}
+
+Value b2p(SPUContext* ctx, const Value& x) { FORCE_DISPATCH(ctx, x); }
+
+Value p2b(SPUContext* ctx, const Value& x) { FORCE_DISPATCH(ctx, x); }
+
+Value a2b(SPUContext* ctx, const Value& x) { TILED_DISPATCH(ctx, x); }
+
+Value b2a(SPUContext* ctx, const Value& x) { TILED_DISPATCH(ctx, x); }
+
+Value and_bp(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value and_bb(SPUContext* ctx, const Value& x, const Value& y) {
+  TILED_DISPATCH(ctx, x, y);
+}
+
+Value xor_bp(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value xor_bb(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value lshift_b(SPUContext* ctx, const Value& x, size_t nbits) {
+  FORCE_DISPATCH(ctx, x, nbits);
+}
+
+Value rshift_b(SPUContext* ctx, const Value& x, size_t nbits) {
+  FORCE_DISPATCH(ctx, x, nbits);
+}
+
+Value arshift_b(SPUContext* ctx, const Value& x, size_t nbits) {
+  FORCE_DISPATCH(ctx, x, nbits);
+}
+
+Value bitrev_b(SPUContext* ctx, const Value& x, size_t start, size_t end) {
+  FORCE_DISPATCH(ctx, x, start, end);
+}
+
+static inline size_t numBits(const Value& in) {
+  return in.storage_type().as<BShare>()->nbits();
+}
+
+static inline Value setNumBits(const Value& in, size_t nbits) {
+  Value out = in;
+  out.storage_type().as<BShare>()->setNbits(nbits);
+  return out;
+}
+
+// TODO: we can not ref api.h, circular reference
+static Value hack_make_p(SPUContext* ctx, uint128_t init,
+                         const std::vector<int64_t>& shape) {
+  return dynDispatch(ctx, "make_p", init, Shape(shape));
+}
+
+Value bitintl_b(SPUContext* ctx, const Value& x, size_t stride) {
+  TRY_DISPATCH(ctx, x, stride);
+
+  // default implementation.
+  // algorithm:
+  //      0000000011111111
+  // swap     ^^^^^^^^
+  //      0000111100001111
+  // swap   ^^^^    ^^^^
+  //      0011001100110011
+  // swap  ^^  ^^  ^^  ^^
+  //      0101010101010101
+  const size_t nbits = x.storage_type().as<BShare>()->nbits();
+  SPU_ENFORCE(absl::has_single_bit(nbits));
+
+  Value out;
+  for (int64_t idx = Log2Ceil(nbits) - 2; idx >= static_cast<int64_t>(stride);
+       idx--) {
+    auto K = hack_make_p(ctx, spu::detail::kBitIntlKeepMasks[idx], x.shape());
+    auto M = hack_make_p(ctx, spu::detail::kBitIntlSwapMasks[idx], x.shape());
+    int64_t S = 1 << idx;
+    // out = (out & K) ^ ((out >> S) & M) ^ ((out & M) << S);
+    out = xor_bb(
+        ctx,
+        xor_bb(ctx, and_bp(ctx, out, K), and_bp(ctx, rshift_b(ctx, out, S), M)),
+        lshift_b(ctx, and_bp(ctx, out, M), S));
+  }
+  out = setNumBits(out, numBits(x));
+  return out;
+}
+
+Value bitdeintl_b(SPUContext* ctx, const Value& x, size_t stride) {
+  TRY_DISPATCH(ctx, x, stride);
+
+  // algorithm:
+  //      0101010101010101
+  // swap  ^^  ^^  ^^  ^^
+  //      0011001100110011
+  // swap   ^^^^    ^^^^
+  //      0000111100001111
+  // swap     ^^^^^^^^
+  //      0000000011111111
+  const size_t nbits = x.storage_type().as<BShare>()->nbits();
+  SPU_ENFORCE(absl::has_single_bit(nbits));
+
+  Value out;
+  for (int64_t idx = stride; idx + 1 < Log2Ceil(nbits); idx++) {
+    auto K = hack_make_p(ctx, spu::detail::kBitIntlKeepMasks[idx], x.shape());
+    auto M = hack_make_p(ctx, spu::detail::kBitIntlSwapMasks[idx], x.shape());
+    int64_t S = 1 << idx;
+    // out = (out & K) ^ ((out >> S) & M) ^ ((out & M) << S);
+    out = xor_bb(
+        ctx,
+        xor_bb(ctx, and_bp(ctx, out, K), and_bp(ctx, rshift_b(ctx, out, S), M)),
+        lshift_b(ctx, and_bp(ctx, out, M), S));
+  }
+  out = setNumBits(out, numBits(x));
+  return out;
+}
+
+namespace {
+
+// TODO: move this to RuntimeConfig
+enum class CircuitType {
+  KoggeStone,
+  Sklansky,
+  Count,
+};
+
+// The kogge-stone adder.
+//
+// P stands for propagate, G stands for generate, where:
+//  (G0, P0) = (g0, p0)
+//  (Gi, Pi) = (gi, pi) o (Gi-1, Pi-1)
+//
+// The `o` here is:
+//  (G0, P0) o (G1, P1) = (G0 ^ (P0 & G1), P0 & P1)
+//
+// Latency log(k) + 1
+Value ppa_kogge_stone(SPUContext* ctx, const Value& lhs, const Value& rhs,
+                      size_t nbits) {
+  // Generate p & g.
+  auto P = xor_bb(ctx, lhs, rhs);
+  auto G = and_bb(ctx, lhs, rhs);
+
+  for (int idx = 0; idx < Log2Ceil(nbits); ++idx) {
+    const size_t offset = 1UL << idx;
+    auto G1 = lshift_b(ctx, G, offset);
+    auto P1 = lshift_b(ctx, P, offset);
+
+    // P1 = P & P1
+    // G1 = G ^ (P & G1)
+    std::vector<Value> res = spu::vectorize(
+        {P, P}, {P1, G1},
+        [&](const Value& xx, const Value& yy) { return and_bb(ctx, xx, yy); });
+    P = std::move(res[0]);
+    G = xor_bb(ctx, G, res[1]);
+  }
+
+  // out = (G << 1) ^ p0
+  auto C = lshift_b(ctx, G, 1);
+  return xor_bb(ctx, xor_bb(ctx, lhs, rhs), C);
+}
+
+std::pair<Value, Value> bit_scatter(SPUContext* ctx, const Value& in,
+                                    size_t stride) {
+  // TODO: use faster bit scatter implementation for ABY3
+  const size_t nbits = numBits(in);
+  SPU_ENFORCE(absl::has_single_bit(nbits), "unsupported {}", nbits);
+  auto out = bitdeintl_b(ctx, in, stride);
+
+  auto hi = rshift_b(ctx, out, nbits / 2);
+  auto mask = hack_make_p(ctx, (static_cast<uint128_t>(1) << (nbits / 2)) - 1,
+                          in.shape());
+  auto lo = and_bp(ctx, out, mask);
+  return std::make_pair(hi, lo);
+}
+
+Value bit_gather(SPUContext* ctx, const Value& hi, const Value& lo,
+                 size_t stride) {
+  const size_t nbits = numBits(hi);
+  SPU_ENFORCE(absl::has_single_bit(nbits), "unsupported {}", nbits);
+  SPU_ENFORCE(nbits == numBits(lo), "nbits mismatch {}, {}", nbits,
+              numBits(lo));
+
+  auto out = xor_bb(ctx, lshift_b(ctx, hi, nbits), lo);
+  return bitintl_b(ctx, out, stride);
+}
+
+// The sklansky adder.
+Value ppa_sklansky(SPUContext* ctx, Value const& lhs, Value const& rhs,
+                   size_t nbits) {
+  SPU_ENFORCE(lhs.numel() == rhs.numel());
+
+  constexpr std::array<uint128_t, 7> kSelMask = {{
+      yacl::MakeUint128(0xFFFFFFFFFFFFFFFF, 0xFFFFFFFFFFFFFFFF),  // invalid
+      yacl::MakeUint128(0xAAAAAAAAAAAAAAAA, 0xAAAAAAAAAAAAAAAA),  // 10101010
+      yacl::MakeUint128(0x8888888888888888, 0x8888888888888888),  // 10001000
+      yacl::MakeUint128(0x8080808080808080, 0x8080808080808080),  // 10000000
+      yacl::MakeUint128(0x8000800080008000, 0x8000800080008000),  // ...
+      yacl::MakeUint128(0x8000000080000000, 0x8000000080000000),  // ...
+      yacl::MakeUint128(0x8000000000000000, 0x8000000000000000),  // ...
+  }};
+
+  // Generate P & G.
+  auto P = xor_bb(ctx, lhs, rhs);
+  auto G = and_bb(ctx, lhs, rhs);
+
+  const size_t bit_width = numBits(lhs);
+  SPU_ENFORCE(bit_width == numBits(rhs), "nbits mismatch {}, {}", bit_width,
+              numBits(rhs));
+  for (int idx = 0; idx < Log2Ceil(nbits); ++idx) {
+    auto [Ph, Pl] = bit_scatter(ctx, P, idx);
+    auto [Gh, Gl] = bit_scatter(ctx, G, idx);
+    // SPU_ENFORCE(numBits(Ph) == bit_width / 2);
+    // SPU_ENFORCE(numBits(Pl) == bit_width / 2);
+    // SPU_ENFORCE(numBits(Gh) == bit_width / 2);
+    // SPU_ENFORCE(numBits(Gl) == bit_width / 2);
+
+    const auto s_mask = hack_make_p(ctx, kSelMask[idx], lhs.shape());
+    auto Gs = and_bp(ctx, Gl, s_mask);
+    auto Ps = and_bp(ctx, Pl, s_mask);
+    for (int j = 0; j < idx; j++) {
+      Gs = xor_bb(ctx, Gs, rshift_b(ctx, Gs, 1 << j));
+      Ps = xor_bb(ctx, Ps, rshift_b(ctx, Ps, 1 << j));
+    }
+    // SPU_ENFORCE(numBits(Ps) == bit_width / 2);
+    // SPU_ENFORCE(numBits(Gs) == bit_width / 2);
+
+    // Ph = Ph & Ps
+    // Gh = Gh ^ (Ph & Gs)
+    std::vector<Value> PG = spu::vectorize(
+        {Ph, Ph}, {Ps, Gs},
+        [&](const Value& xx, const Value& yy) { return and_bb(ctx, xx, yy); });
+    Ph = std::move(PG[0]);
+    Gh = xor_bb(ctx, Gh, PG[1]);
+    // SPU_ENFORCE(numBits(Gh) == numBits(G) / 2);
+    // SPU_ENFORCE(numBits(Ph) == numBits(P) / 2);
+
+    P = bit_gather(ctx, Ph, Pl, idx);
+    G = bit_gather(ctx, Gh, Gl, idx);
+  }
+
+  // out = (G0 << 1) ^ p0
+  auto C = lshift_b(ctx, G, 1);
+  return xor_bb(ctx, xor_bb(ctx, lhs, rhs), C);
+}
+
+}  // namespace
+
+Value add_bb(SPUContext* ctx, const Value& x, const Value& y) {
+  // TRY_DISPATCH
+  if (ctx->hasKernel(__func__)) {
+    SPU_TRACE_MPC_LEAF(ctx, x, y);
+    return tiledDynDispatch(__func__, ctx, x, y);
+  }
+
+  // default implementation
+  SPU_TRACE_MPC_LEAF(ctx, x, y);
+
+  const size_t nbits = numBits(x);
+  SPU_ENFORCE(nbits == numBits(y), "nbits mismatch {}!={}", nbits, numBits(y));
+
+  auto type = CircuitType::KoggeStone;
+
+  switch (type) {
+    case CircuitType::KoggeStone:
+      return ppa_kogge_stone(ctx, x, y, nbits);
+    case CircuitType::Sklansky:
+      return ppa_sklansky(ctx, x, y, nbits);
+    default:
+      SPU_THROW("unknown circuit type {}", static_cast<uint32_t>(type));
+  }
+}
+
+Value carry_out(SPUContext* ctx, const Value& x, const Value& y, size_t k) {
+  // init P & G
+  auto P = xor_bb(ctx, x, y);
+  auto G = and_bb(ctx, x, y);
+
+  // Use kogge stone layout.
+  while (k > 1) {
+    if (k % 2 != 0) {
+      k += 1;
+      P = lshift_b(ctx, P, 1);
+      G = lshift_b(ctx, G, 1);
+    }
+    auto [P1, P0] = bit_scatter(ctx, P, 0);
+    auto [G1, G0] = bit_scatter(ctx, G, 0);
+
+    // Calculate next-level of P, G
+    //   P = P1 & P0
+    //   G = G1 | (P1 & G0)
+    //     = G1 ^ (P1 & G0)
+    std::vector<Value> v = vectorize(
+        {P0, G0}, {P1, P1},
+        [&](const Value& xx, const Value& yy) { return and_bb(ctx, xx, yy); });
+    P = std::move(v[0]);
+    G = xor_bb(ctx, G1, v[1]);
+    k >>= 1;
+  }
+
+  return G;
+}
+
+}  // namespace spu::mpc
diff --git a/libspu/mpc/ab_api.h b/libspu/mpc/ab_api.h
new file mode 100644
index 00000000..78b6cadc
--- /dev/null
+++ b/libspu/mpc/ab_api.h
@@ -0,0 +1,108 @@
+// Copyright 2021 Ant Group Co., Ltd.
+//
+// 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.
+
+#pragma once
+
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
+
+namespace spu::mpc {
+
+Value a2p(SPUContext* ctx, const Value& x);
+Value p2a(SPUContext* ctx, const Value& x);
+Value msb_a2b(SPUContext* ctx, const Value& x);
+
+Value rand_a(SPUContext* ctx, const Shape& shape);
+Value rand_b(SPUContext* ctx, const Shape& shape);
+
+Value not_a(SPUContext* ctx, const Value& x);
+
+Value equal_ap(SPUContext* ctx, const Value& x, const Value& y);
+Value equal_aa(SPUContext* ctx, const Value& x, const Value& y);
+
+Value add_ap(SPUContext* ctx, const Value& x, const Value& y);
+Value add_aa(SPUContext* ctx, const Value& x, const Value&);
+
+Value mul_ap(SPUContext* ctx, const Value& x, const Value& y);
+Value mul_aa(SPUContext* ctx, const Value& x, const Value& y);
+Value mul_a1b(SPUContext* ctx, const Value& x, const Value& y);
+
+Value lshift_a(SPUContext* ctx, const Value& x, size_t nbits);
+Value trunc_a(SPUContext* ctx, const Value& x, size_t nbits);
+
+Value mmul_ap(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k);
+Value mmul_aa(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k);
+
+Type common_type_b(SPUContext* ctx, const Type& a, const Type& b);
+Value cast_type_b(SPUContext* ctx, const Value& a, const Type& to_type);
+
+Value b2p(SPUContext* ctx, const Value& x);
+Value p2b(SPUContext* ctx, const Value& x);
+
+Value a2b(SPUContext* ctx, const Value& x);
+Value b2a(SPUContext* ctx, const Value& x);
+
+Value and_bp(SPUContext* ctx, const Value& x, const Value& y);
+Value and_bb(SPUContext* ctx, const Value& x, const Value& y);
+
+Value xor_bp(SPUContext* ctx, const Value& x, const Value& y);
+Value xor_bb(SPUContext* ctx, const Value& x, const Value& y);
+
+Value lshift_b(SPUContext* ctx, const Value& x, size_t nbits);
+Value rshift_b(SPUContext* ctx, const Value& x, size_t nbits);
+Value arshift_b(SPUContext* ctx, const Value& x, size_t nbits);
+
+// Bit reverse for binary share.
+Value bitrev_b(SPUContext* ctx, const Value& x, size_t start, size_t end);
+
+// TODO: maybe we should add more general PDEP/PEXT instruction.
+
+/// Bit interleave for binary share.
+//
+// Interleave bits of input, so the upper bits of input are in the even
+// positions and lower bits in the odd. Also called Morton Number.
+//
+//   abcdXYZW -> aXbYcZdW
+//
+// stride represent the log shift of the interleaved distance.
+//
+//   00001111 -> 01010101     stride = 0
+//   00001111 -> 00110011     stride = 1
+//   00001111 -> 00001111     stride = 2
+//
+Value bitintl_b(SPUContext* ctx, const Value& x, size_t stride);
+
+// Bit de-interleave for binary share.
+//
+// The reverse bit interleave method, put the even bits at lower half, and odd
+// bits at upper half.
+//
+//   aXbYcZdW -> abcdXYZW
+//
+// stride represent the log shift of the interleaved distance.
+//
+//   01010101 -> 00001111     stride = 0
+//   00110011 -> 00001111     stride = 1
+//   00001111 -> 00001111     stride = 2
+//
+Value bitdeintl_b(SPUContext* ctx, const Value& x, size_t stride);
+
+Value add_bb(SPUContext* ctx, const Value& x, const Value& y);
+
+// compute the k'th bit of x + y
+Value carry_out(SPUContext* ctx, const Value& x, const Value& y, size_t k);
+
+}  // namespace spu::mpc
diff --git a/libspu/mpc/aby3/BUILD.bazel b/libspu/mpc/aby3/BUILD.bazel
index a8c0a162..cb37c5e6 100644
--- a/libspu/mpc/aby3/BUILD.bazel
+++ b/libspu/mpc/aby3/BUILD.bazel
@@ -33,7 +33,6 @@ spu_cc_library(
         ":boolean",
         ":conversion",
         ":value",
-        "//libspu/mpc/common:ab_kernels",
     ],
 )
 
@@ -42,8 +41,8 @@ spu_cc_test(
     srcs = ["protocol_test.cc"],
     deps = [
         ":protocol",
+        "//libspu/mpc:ab_api_test",
         "//libspu/mpc:api_test",
-        "//libspu/mpc/common:ab_api_test",
     ],
 )
 
@@ -56,7 +55,6 @@ spu_cc_library(
         ":type",
         ":value",
         "//libspu/core:trace",
-        "//libspu/mpc/common:ab_api",
         "//libspu/mpc/common:communicator",
         "//libspu/mpc/common:prg_state",
         "//libspu/mpc/utils:circuits",
@@ -71,7 +69,6 @@ spu_cc_library(
         ":type",
         ":value",
         "//libspu/core:platform_utils",
-        "//libspu/mpc/common:ab_api",
         "//libspu/mpc/common:communicator",
         "//libspu/mpc/common:prg_state",
     ],
@@ -85,7 +82,7 @@ spu_cc_library(
         ":ot",
         ":value",
         "//libspu/core:platform_utils",
-        "//libspu/mpc/common:ab_api",
+        "//libspu/mpc:ab_api",
         "//libspu/mpc/common:communicator",
         "//libspu/mpc/common:prg_state",
         "//libspu/mpc/utils:circuits",
diff --git a/libspu/mpc/aby3/arithmetic.cc b/libspu/mpc/aby3/arithmetic.cc
index 89d55c99..d3656997 100644
--- a/libspu/mpc/aby3/arithmetic.cc
+++ b/libspu/mpc/aby3/arithmetic.cc
@@ -20,11 +20,9 @@
 #include "spdlog/spdlog.h"
 
 #include "libspu/core/array_ref.h"
-#include "libspu/core/trace.h"
 #include "libspu/mpc/aby3/ot.h"
 #include "libspu/mpc/aby3/type.h"
 #include "libspu/mpc/aby3/value.h"
-#include "libspu/mpc/common/ab_api.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/common/prg_state.h"
 #include "libspu/mpc/common/pub2k.h"
@@ -116,9 +114,7 @@ std::vector<uint8_t> ring_cast_boolean(const ArrayRef& x) {
 
 }  // namespace
 
-ArrayRef RandA::proc(KernelEvalContext* ctx, size_t size) {
-  SPU_TRACE_MPC_LEAF(ctx, size);
-
+ArrayRef RandA::proc(KernelEvalContext* ctx, size_t size) const {
   auto* prg_state = ctx->getState<PrgState>();
   const auto field = ctx->getState<Z2kState>()->getDefaultField();
 
@@ -149,8 +145,6 @@ ArrayRef RandA::proc(KernelEvalContext* ctx, size_t size) {
 }
 
 ArrayRef A2P::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   auto* comm = ctx->getState<Communicator>();
   const auto field = in.eltype().as<AShrTy>()->field();
 
@@ -179,8 +173,6 @@ ArrayRef A2P::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 }
 
 ArrayRef P2A::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   auto* comm = ctx->getState<Communicator>();
 
   // TODO: we should expect Pub2kTy instead of Ring2k
@@ -225,8 +217,6 @@ ArrayRef P2A::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 }
 
 ArrayRef NotA::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   auto* comm = ctx->getState<Communicator>();
   const auto* in_ty = in.eltype().as<AShrTy>();
   const auto field = in_ty->field();
@@ -261,8 +251,6 @@ ArrayRef NotA::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 ////////////////////////////////////////////////////////////////////
 ArrayRef AddAP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   auto* comm = ctx->getState<Communicator>();
   const auto* lhs_ty = lhs.eltype().as<AShrTy>();
   const auto* rhs_ty = rhs.eltype().as<Pub2kTy>();
@@ -293,8 +281,6 @@ ArrayRef AddAP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef AddAA::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   const auto* lhs_ty = lhs.eltype().as<AShrTy>();
   const auto* rhs_ty = rhs.eltype().as<AShrTy>();
 
@@ -323,8 +309,6 @@ ArrayRef AddAA::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 ////////////////////////////////////////////////////////////////////
 ArrayRef MulAP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   const auto* lhs_ty = lhs.eltype().as<AShrTy>();
   const auto* rhs_ty = rhs.eltype().as<Pub2kTy>();
 
@@ -350,8 +334,6 @@ ArrayRef MulAP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef MulAA::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   const auto field = lhs.eltype().as<Ring2k>()->field();
   auto* comm = ctx->getState<Communicator>();
   auto* prg_state = ctx->getState<PrgState>();
@@ -394,8 +376,6 @@ ArrayRef MulAA::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 // - https://eprint.iacr.org/2018/403.pdf
 ArrayRef MulA1B::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                       const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   SPU_ENFORCE(lhs.numel() == rhs.numel());
   SPU_ENFORCE(lhs.eltype().isa<AShrTy>());
   SPU_ENFORCE(rhs.eltype().isa<BShrTy>() &&
@@ -539,8 +519,6 @@ ArrayRef MulA1B::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 ////////////////////////////////////////////////////////////////////
 ArrayRef MatMulAP::proc(KernelEvalContext* ctx, const ArrayRef& x,
                         const ArrayRef& y, size_t m, size_t n, size_t k) const {
-  SPU_TRACE_MPC_LEAF(ctx, x, y);
-
   const auto field = x.eltype().as<Ring2k>()->field();
 
   ArrayRef z(makeType<AShrTy>(field), m * n);
@@ -559,8 +537,6 @@ ArrayRef MatMulAP::proc(KernelEvalContext* ctx, const ArrayRef& x,
 
 ArrayRef MatMulAA::proc(KernelEvalContext* ctx, const ArrayRef& x,
                         const ArrayRef& y, size_t m, size_t n, size_t k) const {
-  SPU_TRACE_MPC_LEAF(ctx, x, y);
-
   const auto field = x.eltype().as<Ring2k>()->field();
   auto* comm = ctx->getState<Communicator>();
   auto* prg_state = ctx->getState<PrgState>();
@@ -595,8 +571,6 @@ ArrayRef MatMulAA::proc(KernelEvalContext* ctx, const ArrayRef& x,
 
 ArrayRef LShiftA::proc(KernelEvalContext* ctx, const ArrayRef& in,
                        size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, bits);
-
   const auto* in_ty = in.eltype().as<AShrTy>();
   const auto field = in_ty->field();
 
@@ -622,8 +596,6 @@ ArrayRef LShiftA::proc(KernelEvalContext* ctx, const ArrayRef& in,
 // - https://eprint.iacr.org/2018/403.pdf
 ArrayRef TruncA::proc(KernelEvalContext* ctx, const ArrayRef& in,
                       size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, bits);
-
   const auto field = in.eltype().as<Ring2k>()->field();
   auto* prg_state = ctx->getState<PrgState>();
   auto* comm = ctx->getState<Communicator>();
@@ -687,8 +659,6 @@ std::vector<T> openWith(Communicator* comm, size_t peer_rank,
 // - https://arxiv.org/pdf/1910.12435.pdf
 ArrayRef TruncAPr::proc(KernelEvalContext* ctx, const ArrayRef& in,
                         size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, bits);
-
   const auto field = in.eltype().as<AShrTy>()->field();
   const auto numel = in.numel();
   const size_t k = SizeOf(field) * 8;
diff --git a/libspu/mpc/aby3/arithmetic.h b/libspu/mpc/aby3/arithmetic.h
index 28f2d6fe..0dbb7436 100644
--- a/libspu/mpc/aby3/arithmetic.h
+++ b/libspu/mpc/aby3/arithmetic.h
@@ -37,7 +37,7 @@ class A2P : public UnaryKernel {
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in) const override;
 };
 
-class RandA : public Kernel {
+class RandA : public RandKernel {
  public:
   static constexpr char kBindName[] = "rand_a";
 
@@ -45,11 +45,7 @@ class RandA : public Kernel {
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<size_t>(0)));
-  }
-
-  static ArrayRef proc(KernelEvalContext* ctx, size_t size);
+  ArrayRef proc(KernelEvalContext* ctx, size_t size) const override;
 };
 
 class P2A : public UnaryKernel {
diff --git a/libspu/mpc/aby3/boolean.cc b/libspu/mpc/aby3/boolean.cc
index 3c5820f7..eaae0f37 100644
--- a/libspu/mpc/aby3/boolean.cc
+++ b/libspu/mpc/aby3/boolean.cc
@@ -19,7 +19,6 @@
 #include "libspu/core/bit_utils.h"
 #include "libspu/core/parallel_utils.h"
 #include "libspu/core/platform_utils.h"
-#include "libspu/core/trace.h"
 #include "libspu/mpc/aby3/type.h"
 #include "libspu/mpc/aby3/value.h"
 #include "libspu/mpc/common/communicator.h"
@@ -32,8 +31,6 @@ void CommonTypeB::evaluate(KernelEvalContext* ctx) const {
   const Type& lhs = ctx->getParam<Type>(0);
   const Type& rhs = ctx->getParam<Type>(1);
 
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   const size_t lhs_nbits = lhs.as<BShrTy>()->nbits();
   const size_t rhs_nbits = rhs.as<BShrTy>()->nbits();
 
@@ -43,12 +40,8 @@ void CommonTypeB::evaluate(KernelEvalContext* ctx) const {
   ctx->setOutput(makeType<BShrTy>(out_btype, out_nbits));
 }
 
-void CastTypeB::evaluate(KernelEvalContext* ctx) const {
-  const auto& in = ctx->getParam<ArrayRef>(0);
-  const auto& to_type = ctx->getParam<Type>(1);
-
-  SPU_TRACE_MPC_LEAF(ctx, in, to_type);
-
+ArrayRef CastTypeB::proc(KernelEvalContext* ctx, const ArrayRef& in,
+                         const Type& to_type) const {
   ArrayRef out(to_type, in.numel());
   DISPATCH_UINT_PT_TYPES(in.eltype().as<BShrTy>()->getBacktype(), "_", [&]() {
     using InT = ScalarT;
@@ -64,12 +57,10 @@ void CastTypeB::evaluate(KernelEvalContext* ctx) const {
     });
   });
 
-  ctx->setOutput(out);
+  return out;
 }
 
 ArrayRef B2P::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   auto* comm = ctx->getState<Communicator>();
   const PtType btype = in.eltype().as<BShrTy>()->getBacktype();
   const auto field = ctx->getState<Z2kState>()->getDefaultField();
@@ -98,8 +89,6 @@ ArrayRef B2P::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 }
 
 ArrayRef P2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   auto* comm = ctx->getState<Communicator>();
   const auto* in_ty = in.eltype().as<Pub2kTy>();
   const auto field = in_ty->field();
@@ -133,8 +122,6 @@ ArrayRef P2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 
 ArrayRef AndBP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   const auto* lhs_ty = lhs.eltype().as<BShrTy>();
   const auto* rhs_ty = rhs.eltype().as<Pub2kTy>();
 
@@ -167,8 +154,6 @@ ArrayRef AndBP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef AndBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   auto* prg_state = ctx->getState<PrgState>();
   auto* comm = ctx->getState<Communicator>();
 
@@ -216,8 +201,6 @@ ArrayRef AndBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef XorBP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   const auto* lhs_ty = lhs.eltype().as<BShrTy>();
   const auto* rhs_ty = rhs.eltype().as<Pub2kTy>();
 
@@ -250,8 +233,6 @@ ArrayRef XorBP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef XorBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   const auto* lhs_ty = lhs.eltype().as<BShrTy>();
   const auto* rhs_ty = rhs.eltype().as<BShrTy>();
 
@@ -284,8 +265,6 @@ ArrayRef XorBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef LShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
                        size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, bits);
-
   const auto* in_ty = in.eltype().as<BShrTy>();
 
   // TODO: the hal dtype should tell us about the max number of possible bits.
@@ -316,8 +295,6 @@ ArrayRef LShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
 
 ArrayRef RShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
                        size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, bits);
-
   const auto* in_ty = in.eltype().as<BShrTy>();
 
   bits = std::min(in_ty->nbits(), bits);
@@ -348,8 +325,6 @@ ArrayRef RShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
 
 ArrayRef ARShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
                         size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, bits);
-
   const auto field = ctx->getState<Z2kState>()->getDefaultField();
   const auto* in_ty = in.eltype().as<BShrTy>();
 
@@ -378,8 +353,6 @@ ArrayRef ARShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
 
 ArrayRef BitrevB::proc(KernelEvalContext* ctx, const ArrayRef& in, size_t start,
                        size_t end) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, start, end);
-
   SPU_ENFORCE(start <= end && end <= 128);
 
   const auto* in_ty = in.eltype().as<BShrTy>();
@@ -419,12 +392,9 @@ ArrayRef BitrevB::proc(KernelEvalContext* ctx, const ArrayRef& in, size_t start,
   });
 }
 
-void BitIntlB::evaluate(KernelEvalContext* ctx) const {
-  const auto& in = ctx->getParam<ArrayRef>(0);
-  const size_t stride = ctx->getParam<size_t>(1);
-
-  SPU_TRACE_MPC_LEAF(ctx, in, stride);
-
+ArrayRef BitIntlB::proc(KernelEvalContext* ctx, const ArrayRef& in,
+                        size_t stride) const {
+  // void BitIntlB::evaluate(KernelEvalContext* ctx) const {
   const auto* in_ty = in.eltype().as<BShrTy>();
   const size_t nbits = in_ty->nbits();
   SPU_ENFORCE(absl::has_single_bit(nbits));
@@ -441,15 +411,11 @@ void BitIntlB::evaluate(KernelEvalContext* ctx) const {
     });
   });
 
-  ctx->setOutput(out);
+  return out;
 }
 
-void BitDeintlB::evaluate(KernelEvalContext* ctx) const {
-  const auto& in = ctx->getParam<ArrayRef>(0);
-  const size_t stride = ctx->getParam<size_t>(1);
-
-  SPU_TRACE_MPC_LEAF(ctx, in, stride);
-
+ArrayRef BitDeintlB::proc(KernelEvalContext* ctx, const ArrayRef& in,
+                          size_t stride) const {
   const auto* in_ty = in.eltype().as<BShrTy>();
   const size_t nbits = in_ty->nbits();
   SPU_ENFORCE(absl::has_single_bit(nbits));
@@ -466,7 +432,7 @@ void BitDeintlB::evaluate(KernelEvalContext* ctx) const {
     });
   });
 
-  ctx->setOutput(out);
+  return out;
 }
 
 }  // namespace spu::mpc::aby3
diff --git a/libspu/mpc/aby3/boolean.h b/libspu/mpc/aby3/boolean.h
index 33cdb9e3..39ace351 100644
--- a/libspu/mpc/aby3/boolean.h
+++ b/libspu/mpc/aby3/boolean.h
@@ -31,7 +31,7 @@ class CommonTypeB : public Kernel {
   void evaluate(KernelEvalContext* ctx) const override;
 };
 
-class CastTypeB : public Kernel {
+class CastTypeB : public CastTypeKernel {
  public:
   static constexpr char kBindName[] = "cast_type_b";
 
@@ -39,7 +39,8 @@ class CastTypeB : public Kernel {
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override;
+  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
+                const Type& to_type) const override;
 };
 
 class B2P : public UnaryKernel {
@@ -172,7 +173,7 @@ class BitrevB : public BitrevKernel {
                 size_t end) const override;
 };
 
-class BitIntlB : public Kernel {
+class BitIntlB : public BitSplitKernel {
  public:
   static constexpr char kBindName[] = "bitintl_b";
 
@@ -180,10 +181,11 @@ class BitIntlB : public Kernel {
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override;
+  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
+                size_t stride) const override;
 };
 
-class BitDeintlB : public Kernel {
+class BitDeintlB : public BitSplitKernel {
  public:
   static constexpr char kBindName[] = "bitdeintl_b";
 
@@ -191,7 +193,8 @@ class BitDeintlB : public Kernel {
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override;
+  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
+                size_t stride) const override;
 };
 
 }  // namespace spu::mpc::aby3
diff --git a/libspu/mpc/aby3/conversion.cc b/libspu/mpc/aby3/conversion.cc
index b1b8e443..dc42a833 100644
--- a/libspu/mpc/aby3/conversion.cc
+++ b/libspu/mpc/aby3/conversion.cc
@@ -18,16 +18,25 @@
 
 #include "libspu/core/parallel_utils.h"
 #include "libspu/core/platform_utils.h"
-#include "libspu/core/trace.h"
+#include "libspu/core/prelude.h"
+#include "libspu/mpc/ab_api.h"
 #include "libspu/mpc/aby3/type.h"
 #include "libspu/mpc/aby3/value.h"
-#include "libspu/mpc/common/ab_api.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/common/prg_state.h"
 #include "libspu/mpc/common/pub2k.h"
 
 namespace spu::mpc::aby3 {
 
+static ArrayRef wrap_add_bb(SPUContext* ctx, const ArrayRef& x,
+                            const ArrayRef& y) {
+  SPU_ENFORCE(x.numel() == y.numel());
+  const Shape shape = {x.numel()};
+  auto [res, _s, _t] =
+      UnwrapValue(add_bb(ctx, WrapValue(x, shape), WrapValue(y, shape)));
+  return res;
+}
+
 // Reference:
 // ABY3: A Mixed Protocol Framework for Machine Learning
 // P16 5.3 Share Conversions, Bit Decomposition
@@ -35,8 +44,6 @@ namespace spu::mpc::aby3 {
 //
 // Latency: 2 + log(nbits) from 1 rotate and 1 ppa.
 ArrayRef A2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   const auto field = in.eltype().as<Ring2k>()->field();
 
   auto* comm = ctx->getState<Communicator>();
@@ -95,7 +102,7 @@ ArrayRef A2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
     });
   });
 
-  return add_bb(ctx->caller(), m, n);  // comm => log(k) + 1, 2k(logk) + k
+  return wrap_add_bb(ctx->sctx(), m, n);  // comm => log(k) + 1, 2k(logk) + k
 }
 
 ArrayRef B2ASelector::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
@@ -125,8 +132,6 @@ ArrayRef B2ASelector::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 //
 // TODO: convert to single share, will reduce number of rotate.
 ArrayRef B2AByPPA::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   const auto field = ctx->getState<Z2kState>()->getDefaultField();
   const auto* in_ty = in.eltype().as<BShrTy>();
   const size_t in_nbits = in_ty->nbits();
@@ -194,7 +199,7 @@ ArrayRef B2AByPPA::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
       });
 
       // comm => log(k) + 1, 2k(logk) + k
-      auto x_plus_r = add_bb(ctx->caller(), x, r);
+      auto x_plus_r = wrap_add_bb(ctx->sctx(), x, r);
       auto _x_plus_r = ArrayView<std::array<AShrT, 2>>(x_plus_r);
 
       // reveal
@@ -284,8 +289,6 @@ static std::vector<T> bitCompose(absl::Span<T const> in, size_t nbits) {
 // (the receiver) with input bit b2 learns the message c2 (not m[b2]) in the
 // first round, totaling 6k bits and 1 round.
 ArrayRef B2AByOT::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   const auto field = ctx->getState<Z2kState>()->getDefaultField();
   const auto* in_ty = in.eltype().as<BShrTy>();
   const size_t in_nbits = in_ty->nbits();
@@ -424,10 +427,10 @@ ArrayRef B2AByOT::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
   return out;
 }
 
-namespace {
+// TODO: Accelerate bit scatter.
 // split even and odd bits. e.g.
 //   xAyBzCwD -> (xyzw, ABCD)
-std::pair<ArrayRef, ArrayRef> bit_split(const ArrayRef& in) {
+[[maybe_unused]] std::pair<ArrayRef, ArrayRef> bit_split(const ArrayRef& in) {
   constexpr std::array<uint128_t, 6> kSwapMasks = {{
       yacl::MakeUint128(0x2222222222222222, 0x2222222222222222),  // 4bit
       yacl::MakeUint128(0x0C0C0C0C0C0C0C0C, 0x0C0C0C0C0C0C0C0C),  // 8bit
@@ -507,44 +510,7 @@ std::pair<ArrayRef, ArrayRef> bit_split(const ArrayRef& in) {
   return std::make_pair(hi, lo);
 }
 
-// compute the k'th bit of x + y
-ArrayRef carry_out(Object* ctx, const ArrayRef& x, const ArrayRef& y,
-                   size_t k) {
-  // init P & G
-  auto P = xor_bb(ctx, x, y);
-  auto G = and_bb(ctx, x, y);
-
-  // Use kogge stone layout.
-  while (k > 1) {
-    if (k % 2 != 0) {
-      k += 1;
-      P = lshift_b(ctx, P, 1);
-      G = lshift_b(ctx, G, 1);
-    }
-    auto [P1, P0] = bit_split(P);
-    auto [G1, G0] = bit_split(G);
-
-    // Calculate next-level of P, G
-    //   P = P1 & P0
-    //   G = G1 | (P1 & G0)
-    //     = G1 ^ (P1 & G0)
-    std::vector<ArrayRef> v = vectorize(
-        {P0, G0}, {P1, P1}, [&](const ArrayRef& xx, const ArrayRef& yy) {
-          return and_bb(ctx, xx, yy);
-        });
-    P = std::move(v[0]);
-    G = xor_bb(ctx, G1, v[1]);
-    k >>= 1;
-  }
-
-  return G;
-}
-
-}  // namespace
-
 ArrayRef MsbA2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   const auto field = in.eltype().as<AShrTy>()->field();
   const auto numel = in.numel();
   auto* comm = ctx->getState<Communicator>();
@@ -596,12 +562,22 @@ ArrayRef MsbA2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 
   // Compute the k-1'th carry bit.
   size_t nbits = SizeOf(field) * 8 - 1;
-  auto carry = carry_out(ctx->caller(), m, n, nbits);
+  auto* sctx = ctx->sctx();
+
+  const Shape shape = {in.numel()};
+  auto wrap_m = WrapValue(m, shape);
+  auto wrap_n = WrapValue(n, shape);
+  {
+    auto carry = carry_out(sctx, wrap_m, wrap_n, nbits);
 
-  // Compute the k'th bit.
-  //   (m^n)[k] ^ carry
-  auto* obj = ctx->caller();
-  return xor_bb(obj, rshift_b(obj, xor_bb(obj, m, n), nbits), carry);
+    // Compute the k'th bit.
+    //   (m^n)[k] ^ carry
+    auto msb = xor_bb(sctx, rshift_b(sctx, xor_bb(sctx, wrap_m, wrap_n), nbits),
+                      carry);
+
+    auto [msb_data, _shape, _dtype] = UnwrapValue(msb);
+    return msb_data;
+  }
 }
 
 }  // namespace spu::mpc::aby3
diff --git a/libspu/mpc/aby3/protocol.cc b/libspu/mpc/aby3/protocol.cc
index 32c35593..f32f3455 100644
--- a/libspu/mpc/aby3/protocol.cc
+++ b/libspu/mpc/aby3/protocol.cc
@@ -18,81 +18,78 @@
 #include "libspu/mpc/aby3/boolean.h"
 #include "libspu/mpc/aby3/conversion.h"
 #include "libspu/mpc/aby3/type.h"
-#include "libspu/mpc/common/ab_api.h"
-#include "libspu/mpc/common/ab_kernels.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/common/prg_state.h"
 #include "libspu/mpc/common/pub2k.h"
 
 namespace spu::mpc {
 
-std::unique_ptr<Object> makeAby3Protocol(
-    const RuntimeConfig& conf,
-    const std::shared_ptr<yacl::link::Context>& lctx) {
+void regAby3Protocol(SPUContext* ctx,
+                     const std::shared_ptr<yacl::link::Context>& lctx) {
   aby3::registerTypes();
 
-  // FIXME: use same id for different rank
-  auto obj =
-      std::make_unique<Object>(fmt::format("{}-{}", lctx->Rank(), "ABY3"));
-
-  obj->addState<Z2kState>(conf.field());
+  ctx->prot()->addState<Z2kState>(ctx->config().field());
 
   // add communicator
-  obj->addState<Communicator>(lctx);
+  ctx->prot()->addState<Communicator>(lctx);
 
   // register random states & kernels.
-  obj->addState<PrgState>(lctx);
+  ctx->prot()->addState<PrgState>(lctx);
 
   // register public kernels.
-  regPub2kKernels(obj.get());
-
-  // register api kernels
-  regABKernels(obj.get());
+  regPub2kKernels(ctx->prot());
 
   // register arithmetic & binary kernels
-  obj->regKernel<aby3::P2A>();
-  obj->regKernel<aby3::A2P>();
-  obj->regKernel<aby3::NotA>();
-  obj->regKernel<aby3::AddAP>();
-  obj->regKernel<aby3::AddAA>();
-  obj->regKernel<aby3::MulAP>();
-  obj->regKernel<aby3::MulAA>();
-  obj->regKernel<aby3::MulA1B>();
-  obj->regKernel<aby3::MatMulAP>();
-  obj->regKernel<aby3::MatMulAA>();
-  obj->regKernel<aby3::LShiftA>();
+  ctx->prot()->regKernel<aby3::P2A>();
+  ctx->prot()->regKernel<aby3::A2P>();
+  ctx->prot()->regKernel<aby3::NotA>();
+  ctx->prot()->regKernel<aby3::AddAP>();
+  ctx->prot()->regKernel<aby3::AddAA>();
+  ctx->prot()->regKernel<aby3::MulAP>();
+  ctx->prot()->regKernel<aby3::MulAA>();
+  ctx->prot()->regKernel<aby3::MulA1B>();
+  ctx->prot()->regKernel<aby3::MatMulAP>();
+  ctx->prot()->regKernel<aby3::MatMulAA>();
+  ctx->prot()->regKernel<aby3::LShiftA>();
 
 #define ENABLE_PRECISE_ABY3_TRUNCPR
 #ifdef ENABLE_PRECISE_ABY3_TRUNCPR
-  obj->regKernel<aby3::TruncAPr>();
+  ctx->prot()->regKernel<aby3::TruncAPr>();
 #else
-  obj->regKernel<aby3::TruncA>();
+  ctx->prot()->regKernel<aby3::TruncA>();
 #endif
 
-  obj->regKernel<aby3::MsbA2B>();
-
-  obj->regKernel<aby3::CommonTypeB>();
-  obj->regKernel<aby3::CastTypeB>();
-  obj->regKernel<aby3::B2P>();
-  obj->regKernel<aby3::P2B>();
-  obj->regKernel<common::AddBB>();
-  obj->regKernel<aby3::A2B>();
-  obj->regKernel<aby3::B2ASelector>();
-  // obj->regKernel<aby3::B2AByOT>();
-  // obj->regKernel<aby3::B2AByPPA>();
-  obj->regKernel<aby3::AndBP>();
-  obj->regKernel<aby3::AndBB>();
-  obj->regKernel<aby3::XorBP>();
-  obj->regKernel<aby3::XorBB>();
-  obj->regKernel<aby3::LShiftB>();
-  obj->regKernel<aby3::RShiftB>();
-  obj->regKernel<aby3::ARShiftB>();
-  obj->regKernel<aby3::BitrevB>();
-  obj->regKernel<aby3::BitIntlB>();
-  obj->regKernel<aby3::BitDeintlB>();
-  obj->regKernel<aby3::RandA>();
-
-  return obj;
+  ctx->prot()->regKernel<aby3::MsbA2B>();
+
+  ctx->prot()->regKernel<aby3::CommonTypeB>();
+  ctx->prot()->regKernel<aby3::CastTypeB>();
+  ctx->prot()->regKernel<aby3::B2P>();
+  ctx->prot()->regKernel<aby3::P2B>();
+  ctx->prot()->regKernel<aby3::A2B>();
+  ctx->prot()->regKernel<aby3::B2ASelector>();
+  // ctx->prot()->regKernel<aby3::B2AByOT>();
+  // ctx->prot()->regKernel<aby3::B2AByPPA>();
+  ctx->prot()->regKernel<aby3::AndBP>();
+  ctx->prot()->regKernel<aby3::AndBB>();
+  ctx->prot()->regKernel<aby3::XorBP>();
+  ctx->prot()->regKernel<aby3::XorBB>();
+  ctx->prot()->regKernel<aby3::LShiftB>();
+  ctx->prot()->regKernel<aby3::RShiftB>();
+  ctx->prot()->regKernel<aby3::ARShiftB>();
+  ctx->prot()->regKernel<aby3::BitrevB>();
+  ctx->prot()->regKernel<aby3::BitIntlB>();
+  ctx->prot()->regKernel<aby3::BitDeintlB>();
+  ctx->prot()->regKernel<aby3::RandA>();
+}
+
+std::unique_ptr<SPUContext> makeAby3Protocol(
+    const RuntimeConfig& conf,
+    const std::shared_ptr<yacl::link::Context>& lctx) {
+  auto ctx = std::make_unique<SPUContext>(conf, lctx);
+
+  regAby3Protocol(ctx.get(), lctx);
+
+  return ctx;
 }
 
 }  // namespace spu::mpc
diff --git a/libspu/mpc/aby3/protocol.h b/libspu/mpc/aby3/protocol.h
index 474e4f73..8bfe8ae8 100644
--- a/libspu/mpc/aby3/protocol.h
+++ b/libspu/mpc/aby3/protocol.h
@@ -16,12 +16,15 @@
 
 #include "yacl/link/link.h"
 
-#include "libspu/mpc/object.h"
+#include "libspu/core/context.h"
 
 namespace spu::mpc {
 
-std::unique_ptr<Object> makeAby3Protocol(
+std::unique_ptr<SPUContext> makeAby3Protocol(
     const RuntimeConfig& conf,
     const std::shared_ptr<yacl::link::Context>& lctx);
 
+void regAby3Protocol(SPUContext* ctx,
+                     const std::shared_ptr<yacl::link::Context>& lctx);
+
 }  // namespace spu::mpc
diff --git a/libspu/mpc/aby3/type.h b/libspu/mpc/aby3/type.h
index acd02708..f50a17d5 100644
--- a/libspu/mpc/aby3/type.h
+++ b/libspu/mpc/aby3/type.h
@@ -15,7 +15,6 @@
 #pragma once
 
 #include "libspu/core/type.h"
-#include "libspu/mpc/object.h"
 
 namespace spu::mpc::aby3 {
 
diff --git a/libspu/mpc/api.cc b/libspu/mpc/api.cc
index 7d2be256..8ffd8991 100644
--- a/libspu/mpc/api.cc
+++ b/libspu/mpc/api.cc
@@ -1,12 +1,12 @@
 // Copyright 2021 Ant Group Co., Ltd.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
+// you may not use this file except x 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
+// Unless required by applicable law or agreed to x 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
@@ -14,61 +14,377 @@
 
 #include "libspu/mpc/api.h"
 
+#include <future>
+
+#include "libspu/core/trace.h"
+#include "libspu/mpc/ab_api.h"
+
 namespace spu::mpc {
 
-ArrayRef make_p(Object* ctx, uint128_t init, size_t size) {
-  return ctx->call("make_p", init, size);
-}
-
-ArrayRef rand_p(Object* ctx, size_t numel) {
-  return ctx->call("rand_p", numel);
-}
-
-ArrayRef rand_s(Object* ctx, size_t numel) {
-  return ctx->call("rand_s", numel);
-}
-
-Type common_type_s(Object* ctx, const Type& a, const Type& b) {
-  return ctx->call<Type>("common_type_s", a, b);
-}
-
-ArrayRef cast_type_s(Object* ctx, const ArrayRef& a, const Type& to_type) {
-  return ctx->call("cast_type_s", a, to_type);
-}
-
-SPU_MPC_DEF_UNARY_OP(p2s)
-SPU_MPC_DEF_UNARY_OP(s2p)
-SPU_MPC_DEF_UNARY_OP(not_s)
-SPU_MPC_DEF_UNARY_OP(not_p)
-SPU_MPC_DEF_UNARY_OP(msb_s)
-SPU_MPC_DEF_UNARY_OP(msb_p)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(lshift_p)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(lshift_s)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(rshift_p)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(rshift_s)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(arshift_p)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(arshift_s)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(trunc_p)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(trunc_s)
-SPU_MPC_DEF_UNARY_OP_WITH_2SIZE(bitrev_s)
-SPU_MPC_DEF_UNARY_OP_WITH_2SIZE(bitrev_p)
-SPU_MPC_DEF_BINARY_OP(add_pp)
-SPU_MPC_DEF_BINARY_OP(add_sp)
-SPU_MPC_DEF_BINARY_OP(add_ss)
-SPU_MPC_DEF_BINARY_OP(mul_pp)
-SPU_MPC_DEF_BINARY_OP(mul_sp)
-SPU_MPC_DEF_BINARY_OP(mul_ss)
-SPU_MPC_DEF_BINARY_OP(and_pp)
-SPU_MPC_DEF_BINARY_OP(and_sp)
-SPU_MPC_DEF_BINARY_OP(and_ss)
-SPU_MPC_DEF_BINARY_OP(xor_pp)
-SPU_MPC_DEF_BINARY_OP(xor_sp)
-SPU_MPC_DEF_BINARY_OP(xor_ss)
-SPU_MPC_DEF_MMUL(mmul_pp)
-SPU_MPC_DEF_MMUL(mmul_sp)
-SPU_MPC_DEF_MMUL(mmul_ss)
-SPU_MPC_DEF_BINARY_OP(equal_pp);
-SPU_MPC_DEF_OPTIONAL_BINARY_OP(equal_sp);
-SPU_MPC_DEF_OPTIONAL_BINARY_OP(equal_ss);
+#define IsA(x) x.storage_type().isa<AShare>()
+#define IsB(x) x.storage_type().isa<BShare>()
+#define IsP(x) x.storage_type().isa<Public>()
+#define NBits(x) x.storage_type().as<BShare>()->nbits()
+
+// NOLINTBEGIN(readability-identifier-naming)
+Value _2b(SPUContext* ctx, const Value& x) {
+  if (IsA(x)) {
+    return a2b(ctx, x);
+  } else {
+    SPU_ENFORCE(IsB(x), "expect BShare, got {}", x.storage_type());
+    return x;
+  }
+}
+
+Value _2a(SPUContext* ctx, const Value& x) {
+  if (IsB(x)) {
+    return b2a(ctx, x);
+  } else {
+    SPU_ENFORCE(IsA(x), "expect AShare, got {}", x.storage_type());
+    return x;
+  }
+}
+
+// NOLINTEND(readability-identifier-naming)
+
+// TODO: Unify these macros.
+#define FORCE_DISPATCH(CTX, ...)                      \
+  {                                                   \
+    SPU_TRACE_MPC_LEAF(CTX, __VA_ARGS__);             \
+    return dynDispatch((CTX), __func__, __VA_ARGS__); \
+  }
+
+#define TRY_NAMED_DISPATCH(CTX, FNAME, ...)        \
+  if ((CTX)->hasKernel(__func__)) {                \
+    SPU_TRACE_MPC_LEAF(CTX, __VA_ARGS__);          \
+    return dynDispatch((CTX), FNAME, __VA_ARGS__); \
+  }
+
+#define TRY_DISPATCH(CTX, ...) TRY_NAMED_DISPATCH(CTX, __func__, __VA_ARGS__)
+
+Value p2s(SPUContext* ctx, const Value& x) {
+  SPU_TRACE_MPC_DISP(ctx, x);
+
+  TRY_DISPATCH(ctx, x);
+
+  return p2a(ctx, x);
+}
+
+Value s2p(SPUContext* ctx, const Value& x) {
+  SPU_TRACE_MPC_DISP(ctx, x);
+
+  TRY_DISPATCH(ctx, x);
+
+  if (IsA(x)) {
+    return a2p(ctx, x);
+  } else {
+    SPU_ENFORCE(IsB(x));
+    return b2p(ctx, x);
+  }
+}
+
+Value import_s(SPUContext* ctx, const Value& x) {
+  SPU_TRACE_MPC_DISP(ctx, x);
+
+  TRY_DISPATCH(ctx, x);
+
+  SPU_THROW("TODO: import_s not implemented");
+}
+
+Value export_s(SPUContext* ctx, const Value& x, const Type& as_type) {
+  SPU_TRACE_MPC_DISP(ctx, x, as_type);
+
+  TRY_DISPATCH(ctx, x, as_type);
+
+  SPU_THROW("TODO: export_s not implemented");
+}
+
+Type common_type_s(SPUContext* ctx, const Type& a, const Type& b) {
+  SPU_TRACE_MPC_DISP(ctx, a, b);
+
+  // TRY_DISPATCH...
+  if (ctx->hasKernel(__func__)) {
+    SPU_TRACE_MPC_LEAF(ctx, a, b);
+    return dynDispatch<Type>(ctx, __func__, a, b);
+  }
+
+  if (a.isa<AShare>() && b.isa<AShare>()) {
+    SPU_ENFORCE(a == b, "expect same, got a={}, b={}", a, b);
+    return a;
+  } else if (a.isa<AShare>() && b.isa<BShare>()) {
+    return b;
+  } else if (a.isa<BShare>() && b.isa<AShare>()) {
+    return b;
+  } else if (a.isa<BShare>() && b.isa<BShare>()) {
+    return common_type_b(ctx, b, b);
+  } else {
+    SPU_THROW("should not be here, a={}, b={}", a, b);
+  }
+}
+
+Value cast_type_s(SPUContext* ctx, const Value& frm, const Type& to_type) {
+  SPU_TRACE_MPC_DISP(ctx, frm, to_type);
+
+  TRY_DISPATCH(ctx, frm, to_type);
+
+  if (IsA(frm) && to_type.isa<AShare>()) {
+    SPU_ENFORCE(frm.storage_type() == to_type,
+                "expect same, got frm={}, to_type={}", frm, to_type);
+    // do nothing.
+    return frm;
+  } else if (IsA(frm) && to_type.isa<BShare>()) {
+    return a2b(ctx, frm);
+  } else if (IsB(frm) && to_type.isa<AShare>()) {
+    return b2a(ctx, frm);
+  } else if (IsB(frm) && to_type.isa<BShare>()) {
+    return cast_type_b(ctx, frm, to_type);
+  } else {
+    SPU_THROW("should not be here, frm={}, to_type={}", frm, to_type);
+  }
+}
+
+Value make_p(SPUContext* ctx, uint128_t init, const Shape& shape) {
+  FORCE_DISPATCH(ctx, init, shape);
+}
+
+Value rand_p(SPUContext* ctx, const Shape& shape) {
+  FORCE_DISPATCH(ctx, shape);
+}
+
+Value rand_s(SPUContext* ctx, const Shape& shape) {
+  SPU_TRACE_MPC_DISP(ctx, shape);
+  TRY_DISPATCH(ctx, shape);
+  // always return random a share
+  return rand_a(ctx, shape);
+}
+
+Value not_p(SPUContext* ctx, const Value& x) { FORCE_DISPATCH(ctx, x); }
+
+Value not_s(SPUContext* ctx, const Value& x) {
+  SPU_TRACE_MPC_DISP(ctx, x);
+  TRY_DISPATCH(ctx, x);
+  // TODO: Both A&B could handle not(invert).
+  // if (x.eltype().isa<BShare>()) {
+  //  return not_b(ctx, x);
+  //} else {
+  //  SPU_ENFORCE(x.eltype().isa<AShare>());
+  //  return not_a(ctx, x);
+  //}
+  return not_a(ctx, _2a(ctx, x));
+}
+
+Value msb_p(SPUContext* ctx, const Value& x) { FORCE_DISPATCH(ctx, x); }
+
+Value msb_s(SPUContext* ctx, const Value& x) {
+  SPU_TRACE_MPC_DISP(ctx, x);
+  TRY_DISPATCH(ctx, x);
+
+  // TODO: this is buggy.
+  const auto field = x.storage_type().as<Ring2k>()->field();
+
+  if (ctx->hasKernel("msb_a2b")) {
+    if (IsB(x)) {
+      return rshift_b(ctx, x, SizeOf(field) * 8 - 1);
+    } else {
+      // fast path, directly apply msb x AShare, result a BShare.
+      return msb_a2b(ctx, x);
+    }
+  } else {
+    return rshift_b(ctx, _2b(ctx, x), SizeOf(field) * 8 - 1);
+  }
+}
+
+Value equal_pp(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+std::optional<Value> equal_sp(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+
+  if (IsA(x) && ctx->hasKernel("equal_ap")) {
+    return dynDispatch(ctx, "equal_ap", x, y);
+  } else if (IsB(x) && ctx->hasKernel("equal_bp")) {
+    return dynDispatch(ctx, "equal_bp", x, y);
+  }
+
+  return std::nullopt;
+}
+
+std::optional<Value> equal_ss(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+
+  // try fast path
+  // TODO: use cost model instead of hand-coded priority.
+  if (IsA(x) && IsA(y) && ctx->hasKernel("equal_aa")) {
+    return dynDispatch(ctx, "equal_aa", x, y);
+  } else if (IsB(x) && IsB(y) && ctx->hasKernel("equal_bb")) {
+    return dynDispatch(ctx, "equal_bb", x, y);
+  } else if ((IsA(x) && IsB(y)) || (IsB(x) && IsA(y))) {
+    // mixed a & b, both OK, hardcode to a.
+    if (ctx->hasKernel("equal_aa")) {
+      return dynDispatch(ctx, "equal_aa", _2a(ctx, x), _2a(ctx, y));
+    }
+    if (ctx->hasKernel("equal_bb")) {
+      return dynDispatch(ctx, "equal_bb", _2b(ctx, x), _2b(ctx, y));
+    }
+  }
+
+  return std::nullopt;
+}
+
+Value add_pp(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value add_sp(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+  return add_ap(ctx, _2a(ctx, x), y);
+}
+
+Value add_ss(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+  return add_aa(ctx, _2a(ctx, x), _2a(ctx, y));
+}
+
+Value mul_pp(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value mul_sp(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+  return mul_ap(ctx, _2a(ctx, x), y);
+}
+
+static bool hasMulA1B(SPUContext* ctx) { return ctx->hasKernel("mul_a1b"); }
+
+Value mul_ss(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+
+  if (hasMulA1B(ctx) && IsA(y) && IsB(x) && NBits(x) == 1) {
+    return mul_a1b(ctx, y, x);
+  }
+  if (hasMulA1B(ctx) && IsA(x) && IsB(y) && NBits(y) == 1) {
+    return mul_a1b(ctx, x, y);
+  }
+
+  // NOTE(juhou): Multiplication of two bits
+  if (IsB(x) && NBits(x) == 1 && IsB(y) && NBits(y) == 1) {
+    return and_bb(ctx, x, y);
+  }
+  return mul_aa(ctx, _2a(ctx, x), _2a(ctx, y));
+}
+
+Value mmul_pp(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k) {
+  FORCE_DISPATCH(ctx, x, y, m, n, k);
+}
+
+Value mmul_sp(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k) {
+  SPU_TRACE_MPC_DISP(ctx, x, y, m, n, k);
+  TRY_DISPATCH(ctx, x, y, m, n, k);
+  return mmul_ap(ctx, _2a(ctx, x), y, m, n, k);
+}
+
+Value mmul_ss(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y, m, n, k);
+  return mmul_aa(ctx, _2a(ctx, x), _2a(ctx, y), m, n, k);
+}
+
+Value and_pp(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value and_sp(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+  return and_bp(ctx, _2b(ctx, x), y);
+}
+
+Value and_ss(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+  return and_bb(ctx, _2b(ctx, x), _2b(ctx, y));
+}
+
+Value xor_pp(SPUContext* ctx, const Value& x, const Value& y) {
+  FORCE_DISPATCH(ctx, x, y);
+}
+
+Value xor_sp(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+  return xor_bp(ctx, _2b(ctx, x), y);
+}
+
+Value xor_ss(SPUContext* ctx, const Value& x, const Value& y) {
+  SPU_TRACE_MPC_DISP(ctx, x, y);
+  TRY_DISPATCH(ctx, x, y);
+  return xor_bb(ctx, _2b(ctx, x), _2b(ctx, y));
+}
+
+Value lshift_p(SPUContext* ctx, const Value& x, size_t nbits) {
+  FORCE_DISPATCH(ctx, x, nbits);
+}
+
+Value lshift_s(SPUContext* ctx, const Value& x, size_t bits) {
+  SPU_TRACE_MPC_DISP(ctx, x, bits);
+  TRY_DISPATCH(ctx, x, bits);
+  if (IsA(x)) {
+    return lshift_a(ctx, x, bits);
+  } else if (IsB(x)) {
+    return lshift_b(ctx, x, bits);
+  } else {
+    SPU_THROW("Unsupported type {}", x.storage_type());
+  }
+}
+
+Value rshift_p(SPUContext* ctx, const Value& x, size_t nbits) {
+  FORCE_DISPATCH(ctx, x, nbits);
+}
+
+Value rshift_s(SPUContext* ctx, const Value& x, size_t bits) {
+  SPU_TRACE_MPC_DISP(ctx, x, bits);
+  TRY_DISPATCH(ctx, x, bits);
+  return rshift_b(ctx, _2b(ctx, x), bits);
+}
+
+Value arshift_p(SPUContext* ctx, const Value& x, size_t nbits) {
+  FORCE_DISPATCH(ctx, x, nbits);
+}
+
+Value arshift_s(SPUContext* ctx, const Value& x, size_t bits) {
+  SPU_TRACE_MPC_DISP(ctx, x, bits);
+  TRY_DISPATCH(ctx, x, bits);
+  return arshift_b(ctx, _2b(ctx, x), bits);
+}
+
+Value trunc_p(SPUContext* ctx, const Value& x, size_t nbits) {
+  FORCE_DISPATCH(ctx, x, nbits);
+}
+
+Value trunc_s(SPUContext* ctx, const Value& x, size_t bits) {
+  SPU_TRACE_MPC_DISP(ctx, x, bits);
+  TRY_DISPATCH(ctx, x, bits);
+  return trunc_a(ctx, _2a(ctx, x), bits);
+}
+
+Value bitrev_p(SPUContext* ctx, const Value& x, size_t start, size_t end) {
+  FORCE_DISPATCH(ctx, x, start, end);
+}
+
+Value bitrev_s(SPUContext* ctx, const Value& x, size_t start, size_t end) {
+  SPU_TRACE_MPC_DISP(ctx, x, start, end);
+  TRY_DISPATCH(ctx, x, start, end);
+  return bitrev_b(ctx, _2b(ctx, x), start, end);
+}
 
 }  // namespace spu::mpc
diff --git a/libspu/mpc/api.h b/libspu/mpc/api.h
index 86f1f0c4..e1e48453 100644
--- a/libspu/mpc/api.h
+++ b/libspu/mpc/api.h
@@ -16,23 +16,28 @@
 
 #include <optional>
 
-#include "libspu/core/array_ref.h"
-#include "libspu/mpc/object.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 
 namespace spu::mpc {
 
+// TODO: add to naming conventions.
+// - use x,y,z for value
+// - use a,b,c for type
+// - follow current module style.
+
 // Convert a public to a secret.
 //
 // In most of cases, you should not do this, becasue:
 // 1. This only convert the 'type' to secret, but partipants still knows its
 //    value at the moment.
 // 2. Nearly all ops has public parameter overload, we should use it directly.
-ArrayRef p2s(Object* ctx, const ArrayRef&);
+Value p2s(SPUContext* ctx, const Value& x);
 
 // Convert a secret to a public, aka, reveal.
 //
 // Note: this API indicates information leak.
-ArrayRef s2p(Object* ctx, const ArrayRef&);
+Value s2p(SPUContext* ctx, const Value& x);
 
 // Import will be called on all parameters at the beginning program.
 //
@@ -43,7 +48,7 @@ ArrayRef s2p(Object* ctx, const ArrayRef&);
 // @param ctx, the evaluation context.
 // @param in, the type may not be of current protocol's type, but
 //            it should be a Secret type.
-ArrayRef import_s(Object* ctx, const ArrayRef& in);
+Value import_s(SPUContext* ctx, const Value& in);
 
 // Export a secret value as a given type.
 //
@@ -54,7 +59,7 @@ ArrayRef import_s(Object* ctx, const ArrayRef& in);
 // @param ctx, the evaluation context.
 // @param in, the input should be one of current protocol's type.
 // @param as_type, the target type, it should be a Secret type.
-ArrayRef export_s(Object* ctx, const ArrayRef& in, const Type& as_type);
+Value export_s(SPUContext* ctx, const Value& in, const Type& as_type);
 
 // Get the common type of secrets.
 //
@@ -62,101 +67,66 @@ ArrayRef export_s(Object* ctx, const ArrayRef& in, const Type& as_type);
 // formats, like AShare/BShare, which make them not concatable.
 //
 // This api calculate the common type.
-Type common_type_s(Object* ctx, const Type& a, const Type& b);
-ArrayRef cast_type_s(Object* ctx, const ArrayRef& a, const Type& to_type);
+Type common_type_s(SPUContext* ctx, const Type& a, const Type& b);
+Value cast_type_s(SPUContext* ctx, const Value& frm, const Type& to_type);
 
 // Make a public variable with given plaintext input.
 //
 // All parties knowns the value.
-ArrayRef make_p(Object* ctx, uint128_t init, size_t size);
+Value make_p(SPUContext* ctx, uint128_t init, const Shape& shape);
 
 // parties random a public together.
-ArrayRef rand_p(Object* ctx, size_t);
-ArrayRef rand_s(Object* ctx, size_t);
+Value rand_p(SPUContext* ctx, const Shape& shape);
+Value rand_s(SPUContext* ctx, const Shape& shape);
 
 // Compute bitwise_not(invert) of a value in ring 2k space.
-ArrayRef not_s(Object* ctx, const ArrayRef&);
-ArrayRef not_p(Object* ctx, const ArrayRef&);
+Value not_p(SPUContext* ctx, const Value& x);
+Value not_s(SPUContext* ctx, const Value& x);
 
-ArrayRef msb_p(Object* ctx, const ArrayRef&);
-ArrayRef msb_s(Object* ctx, const ArrayRef&);
+Value msb_p(SPUContext* ctx, const Value& x);
+Value msb_s(SPUContext* ctx, const Value& x);
 
-ArrayRef equal_pp(Object* ctx, const ArrayRef&, const ArrayRef&);
+Value equal_pp(SPUContext* ctx, const Value& x, const Value& y);
 // Optional API, return nullopt if no valid kernel found.
-std::optional<ArrayRef> equal_sp(Object* ctx, const ArrayRef&, const ArrayRef&);
-std::optional<ArrayRef> equal_ss(Object* ctx, const ArrayRef&, const ArrayRef&);
+std::optional<Value> equal_sp(SPUContext* ctx, const Value& x, const Value& y);
+std::optional<Value> equal_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-ArrayRef lshift_p(Object* ctx, const ArrayRef&, size_t);
-ArrayRef lshift_s(Object* ctx, const ArrayRef&, size_t);
+Value add_pp(SPUContext* ctx, const Value& x, const Value& y);
+Value add_sp(SPUContext* ctx, const Value& x, const Value& y);
+Value add_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-ArrayRef rshift_p(Object* ctx, const ArrayRef&, size_t);
-ArrayRef rshift_s(Object* ctx, const ArrayRef&, size_t);
+Value mul_pp(SPUContext* ctx, const Value& x, const Value& y);
+Value mul_sp(SPUContext* ctx, const Value& x, const Value& y);
+Value mul_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-ArrayRef arshift_p(Object* ctx, const ArrayRef&, size_t);
-ArrayRef arshift_s(Object* ctx, const ArrayRef&, size_t);
-ArrayRef trunc_p(Object* ctx, const ArrayRef&, size_t);
-ArrayRef trunc_s(Object* ctx, const ArrayRef&, size_t);
+Value and_pp(SPUContext* ctx, const Value& x, const Value& y);
+Value and_sp(SPUContext* ctx, const Value& x, const Value& y);
+Value and_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-// Reverse bit, like MISP BITREV instruction, and linux bitrev library.
-ArrayRef bitrev_s(Object* ctx, const ArrayRef&, size_t, size_t);
-ArrayRef bitrev_p(Object* ctx, const ArrayRef&, size_t, size_t);
+Value xor_pp(SPUContext* ctx, const Value& x, const Value& y);
+Value xor_sp(SPUContext* ctx, const Value& x, const Value& y);
+Value xor_ss(SPUContext* ctx, const Value& x, const Value& y);
 
-ArrayRef add_pp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef add_sp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef add_ss(Object* ctx, const ArrayRef&, const ArrayRef&);
+Value mmul_pp(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k);
+Value mmul_sp(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k);
+Value mmul_ss(SPUContext* ctx, const Value& x, const Value& y, size_t m,
+              size_t n, size_t k);
 
-ArrayRef mul_pp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef mul_sp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef mul_ss(Object* ctx, const ArrayRef&, const ArrayRef&);
+Value lshift_p(SPUContext* ctx, const Value& x, size_t nbits);
+Value lshift_s(SPUContext* ctx, const Value& x, size_t nbits);
 
-ArrayRef and_pp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef and_sp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef and_ss(Object* ctx, const ArrayRef&, const ArrayRef&);
+Value rshift_p(SPUContext* ctx, const Value& x, size_t nbits);
+Value rshift_s(SPUContext* ctx, const Value& x, size_t nbits);
 
-ArrayRef xor_pp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef xor_sp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef xor_ss(Object* ctx, const ArrayRef&, const ArrayRef&);
+Value arshift_p(SPUContext* ctx, const Value& x, size_t nbits);
+Value arshift_s(SPUContext* ctx, const Value& x, size_t nbits);
+Value trunc_p(SPUContext* ctx, const Value& x, size_t nbits);
+Value trunc_s(SPUContext* ctx, const Value& x, size_t nbits);
 
-ArrayRef mmul_pp(Object* ctx, const ArrayRef&, const ArrayRef&, size_t, size_t,
-                 size_t);
-ArrayRef mmul_sp(Object* ctx, const ArrayRef&, const ArrayRef&, size_t, size_t,
-                 size_t);
-ArrayRef mmul_ss(Object* ctx, const ArrayRef&, const ArrayRef&, size_t, size_t,
-                 size_t);
+// Reverse bit, like MISP BITREV instruction, and linux bitrev library.
+Value bitrev_s(SPUContext* ctx, const Value& x, size_t start, size_t end);
+Value bitrev_p(SPUContext* ctx, const Value& x, size_t start, size_t end);
 
 }  // namespace spu::mpc
-
-#define SPU_MPC_DEF_UNARY_OP(NAME)                 \
-  ArrayRef NAME(Object* ctx, const ArrayRef& in) { \
-    return ctx->call(#NAME, in);                   \
-  }
-
-#define SPU_MPC_DEF_UNARY_OP_WITH_SIZE(NAME)                  \
-  ArrayRef NAME(Object* ctx, const ArrayRef& in, size_t sz) { \
-    return ctx->call(#NAME, in, sz);                          \
-  }
-
-#define SPU_MPC_DEF_UNARY_OP_WITH_2SIZE(NAME)                              \
-  ArrayRef NAME(Object* ctx, const ArrayRef& in, size_t sz1, size_t sz2) { \
-    return ctx->call(#NAME, in, sz1, sz2);                                 \
-  }
-
-#define SPU_MPC_DEF_BINARY_OP(NAME)                                  \
-  ArrayRef NAME(Object* ctx, const ArrayRef& x, const ArrayRef& y) { \
-    return ctx->call(#NAME, x, y);                                   \
-  }
-
-#define SPU_MPC_DEF_OPTIONAL_BINARY_OP(NAME)                   \
-  std::optional<ArrayRef> NAME(Object* ctx, const ArrayRef& x, \
-                               const ArrayRef& y) {            \
-    if (!ctx->hasKernel(#NAME)) {                              \
-      return std::nullopt;                                     \
-    }                                                          \
-    return ctx->call<std::optional<ArrayRef>>(#NAME, x, y);    \
-  }
-
-#define SPU_MPC_DEF_MMUL(NAME)                                                 \
-  ArrayRef NAME(Object* ctx, const ArrayRef& x, const ArrayRef& y, size_t sz1, \
-                size_t sz2, size_t sz3) {                                      \
-    return ctx->call(#NAME, x, y, sz1, sz2, sz3);                              \
-  }
diff --git a/libspu/mpc/cheetah/BUILD.bazel b/libspu/mpc/cheetah/BUILD.bazel
index c4fed037..842bc03d 100644
--- a/libspu/mpc/cheetah/BUILD.bazel
+++ b/libspu/mpc/cheetah/BUILD.bazel
@@ -84,8 +84,6 @@ spu_cc_library(
         ":boolean",
         ":conversion",
         ":state",
-        "//libspu/mpc/common:ab_api",
-        "//libspu/mpc/common:ab_kernels",
         "//libspu/mpc/common:prg_state",
         "//libspu/mpc/common:pub2k",
     ],
@@ -98,8 +96,8 @@ spu_cc_test(
     srcs = ["protocol_test.cc"],
     deps = [
         ":protocol",
+        "//libspu/mpc:ab_api_test",
         "//libspu/mpc:api_test",
-        "//libspu/mpc/common:ab_api_test",
     ],
 )
 
diff --git a/libspu/mpc/cheetah/arithmetic.cc b/libspu/mpc/cheetah/arithmetic.cc
index 6a84e104..3d4d945a 100644
--- a/libspu/mpc/cheetah/arithmetic.cc
+++ b/libspu/mpc/cheetah/arithmetic.cc
@@ -69,7 +69,6 @@ ArrayRef TruncAWithSign::proc(KernelEvalContext* ctx, const ArrayRef& x,
 
 ArrayRef TruncA::proc(KernelEvalContext* ctx, const ArrayRef& x,
                       size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, x);
   auto* comm = ctx->getState<Communicator>();
   auto* ot_state = ctx->getState<CheetahOTState>();
   size_t n = x.numel();
@@ -103,8 +102,6 @@ ArrayRef TruncA::proc(KernelEvalContext* ctx, const ArrayRef& x,
 }
 
 ArrayRef MsbA2B::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
-  SPU_TRACE_MPC_LEAF(ctx, x);
-
   auto* comm = ctx->getState<Communicator>();
   auto* ot_state = ctx->getState<CheetahOTState>();
   size_t n = x.numel();
@@ -228,7 +225,6 @@ ArrayRef EqualAA::proc(KernelEvalContext* ctx, const ArrayRef& x,
 
 ArrayRef MulA1B::proc(KernelEvalContext* ctx, const ArrayRef& x,
                       const ArrayRef& y) const {
-  SPU_TRACE_MPC_LEAF(ctx, x, y);
   SPU_ENFORCE_EQ(x.numel(), y.numel());
   auto* comm = ctx->getState<Communicator>();
   auto* ot_state = ctx->getState<CheetahOTState>();
@@ -264,7 +260,6 @@ ArrayRef MulA1B::proc(KernelEvalContext* ctx, const ArrayRef& x,
 
 ArrayRef MulAA::proc(KernelEvalContext* ctx, const ArrayRef& x,
                      const ArrayRef& y) const {
-  SPU_TRACE_MPC_LEAF(ctx, x, y);
   SPU_ENFORCE_EQ(x.numel(), y.numel());
 
   size_t batch_sze = ctx->getState<CheetahMulState>()->get()->OLEBatchSize();
@@ -287,7 +282,7 @@ ArrayRef MulAA::mulWithBeaver(KernelEvalContext* ctx, const ArrayRef& x,
       ctx->getState<CheetahMulState>()->TakeCachedBeaver(field, numel);
   YACL_ENFORCE_EQ(a.numel(), numel);
 
-  auto* comm = ctx->caller()->getState<Communicator>();
+  auto* comm = ctx->getState<Communicator>();
   // Open x - a & y - b
   auto res =
       vectorize({ring_sub(x, a), ring_sub(y, b)}, [&](const ArrayRef& s) {
@@ -338,7 +333,6 @@ ArrayRef MulAA::mulDirectly(KernelEvalContext* ctx, const ArrayRef& x,
 // A is (M, K); B is (K, N)
 ArrayRef MatMulAA::proc(KernelEvalContext* ctx, const ArrayRef& x,
                         const ArrayRef& y, size_t m, size_t n, size_t k) const {
-  SPU_TRACE_MPC_LEAF(ctx, x, y);
   if (0 == x.numel() || 0 == y.numel()) {
     return ArrayRef(x.eltype(), 0);
   }
diff --git a/libspu/mpc/cheetah/arithmetic.h b/libspu/mpc/cheetah/arithmetic.h
index 7d714bb6..5b9ad53e 100644
--- a/libspu/mpc/cheetah/arithmetic.h
+++ b/libspu/mpc/cheetah/arithmetic.h
@@ -16,12 +16,9 @@
 
 #include "libspu/mpc/kernel.h"
 #include "libspu/mpc/semi2k/arithmetic.h"
-#include "libspu/mpc/utils/cexpr.h"
 
 namespace spu::mpc::cheetah {
 
-using ZeroA = spu::mpc::semi2k::ZeroA;
-
 using RandA = spu::mpc::semi2k::RandA;
 
 using P2A = spu::mpc::semi2k::P2A;
diff --git a/libspu/mpc/cheetah/boolean.cc b/libspu/mpc/cheetah/boolean.cc
index dc037067..c4ccc655 100644
--- a/libspu/mpc/cheetah/boolean.cc
+++ b/libspu/mpc/cheetah/boolean.cc
@@ -14,7 +14,6 @@
 
 #include "libspu/mpc/cheetah/boolean.h"
 
-#include "libspu/core/trace.h"
 #include "libspu/mpc/cheetah/state.h"
 #include "libspu/mpc/common/pub2k.h"
 #include "libspu/mpc/utils/ring_ops.h"
@@ -24,7 +23,6 @@ constexpr size_t kMinWorkSize = 5000;
 
 ArrayRef AndBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
   SPU_ENFORCE_EQ(lhs.numel(), rhs.numel());
 
   auto* comm = ctx->getState<Communicator>();
diff --git a/libspu/mpc/cheetah/boolean.h b/libspu/mpc/cheetah/boolean.h
index 67daf791..4aa5a04c 100644
--- a/libspu/mpc/cheetah/boolean.h
+++ b/libspu/mpc/cheetah/boolean.h
@@ -16,7 +16,6 @@
 
 #include "libspu/mpc/kernel.h"
 #include "libspu/mpc/semi2k/boolean.h"
-#include "libspu/mpc/utils/cexpr.h"
 
 namespace spu::mpc::cheetah {
 
@@ -24,8 +23,6 @@ using CommonTypeB = spu::mpc::semi2k::CommonTypeB;
 
 using CastTypeB = spu::mpc::semi2k::CastTypeB;
 
-using ZeroB = spu::mpc::semi2k::ZeroB;
-
 using B2P = spu::mpc::semi2k::B2P;
 
 using P2B = spu::mpc::semi2k::P2B;
diff --git a/libspu/mpc/cheetah/conversion.cc b/libspu/mpc/cheetah/conversion.cc
index d7032755..2046b47f 100644
--- a/libspu/mpc/cheetah/conversion.cc
+++ b/libspu/mpc/cheetah/conversion.cc
@@ -14,7 +14,6 @@
 
 #include "libspu/mpc/cheetah/conversion.h"
 
-#include "libspu/core/trace.h"
 #include "libspu/mpc/cheetah/state.h"
 #include "libspu/mpc/common/pub2k.h"
 #include "libspu/mpc/semi2k/type.h"  // TODO: use cheetah type
@@ -23,7 +22,6 @@ namespace spu::mpc::cheetah {
 constexpr size_t kMinWorkSize = 5000;
 
 ArrayRef B2A::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
-  SPU_TRACE_MPC_LEAF(ctx, x);
   auto* comm = ctx->getState<Communicator>();
   auto* ot_state = ctx->getState<CheetahOTState>();
   size_t n = x.numel();
diff --git a/libspu/mpc/cheetah/conversion.h b/libspu/mpc/cheetah/conversion.h
index a57bb86c..92272c53 100644
--- a/libspu/mpc/cheetah/conversion.h
+++ b/libspu/mpc/cheetah/conversion.h
@@ -16,7 +16,6 @@
 
 #include "libspu/mpc/kernel.h"
 #include "libspu/mpc/semi2k/conversion.h"
-#include "libspu/mpc/utils/cexpr.h"
 
 namespace spu::mpc::cheetah {
 
diff --git a/libspu/mpc/cheetah/protocol.cc b/libspu/mpc/cheetah/protocol.cc
index 869c9a5a..aacfa79f 100644
--- a/libspu/mpc/cheetah/protocol.cc
+++ b/libspu/mpc/cheetah/protocol.cc
@@ -22,86 +22,80 @@
 #include "libspu/mpc/cheetah/boolean.h"
 #include "libspu/mpc/cheetah/conversion.h"
 #include "libspu/mpc/cheetah/state.h"
-#include "libspu/mpc/common/ab_api.h"
-#include "libspu/mpc/common/ab_kernels.h"
 #include "libspu/mpc/common/pub2k.h"
-#include "libspu/mpc/object.h"
 #include "libspu/mpc/semi2k/state.h"
 #include "libspu/mpc/semi2k/type.h"
 
 namespace spu::mpc {
 
-std::unique_ptr<Object> makeCheetahProtocol(
-    const RuntimeConfig& conf,
-    const std::shared_ptr<yacl::link::Context>& lctx) {
+void regCheetahProtocol(SPUContext* ctx,
+                        const std::shared_ptr<yacl::link::Context>& lctx) {
   semi2k::registerTypes();
 
-  // FIXME: use same id for different rank
-  auto obj =
-      std::make_unique<Object>(fmt::format("{}-{}", lctx->Rank(), "CHEETAH"));
-
   // add communicator
-  obj->addState<Communicator>(lctx);
+  ctx->prot()->addState<Communicator>(lctx);
 
   // register random states & kernels.
-  obj->addState<PrgState>(lctx);
+  ctx->prot()->addState<PrgState>(lctx);
 
   // add Z2k state.
-  obj->addState<Z2kState>(conf.field());
+  ctx->prot()->addState<Z2kState>(ctx->config().field());
 
   // add Cheetah states
-  obj->addState<cheetah::CheetahMulState>(lctx);
-  obj->addState<cheetah::CheetahDotState>(lctx);
-  obj->addState<cheetah::CheetahOTState>();
+  ctx->prot()->addState<cheetah::CheetahMulState>(lctx);
+  ctx->prot()->addState<cheetah::CheetahDotState>(lctx);
+  ctx->prot()->addState<cheetah::CheetahOTState>();
 
   // register public kernels.
-  regPub2kKernels(obj.get());
-
-  // register compute kernels
-  regABKernels(obj.get());
+  regPub2kKernels(ctx->prot());
 
   // register arithmetic & binary kernels
-  obj->regKernel<cheetah::ZeroA>();
-  obj->regKernel<cheetah::P2A>();
-  obj->regKernel<cheetah::A2P>();
-  obj->regKernel<cheetah::NotA>();
-  obj->regKernel<cheetah::AddAP>();
-  obj->regKernel<cheetah::AddAA>();
-  obj->regKernel<cheetah::MulAP>();
-  obj->regKernel<cheetah::MulAA>();
-  obj->regKernel<cheetah::MulA1B>();
-  obj->regKernel<cheetah::EqualAA>();
-  obj->regKernel<cheetah::EqualAP>();
-  obj->regKernel<cheetah::MatMulAP>();
-  obj->regKernel<cheetah::MatMulAA>();
-  obj->regKernel<cheetah::Conv2DAA>();
-  obj->regKernel<cheetah::LShiftA>();
-  obj->regKernel<cheetah::TruncA>();
-  obj->regKernel<cheetah::TruncAWithSign>();
-  obj->regKernel<cheetah::MsbA2B>();
-
-  obj->regKernel<common::AddBB>();
-  obj->regKernel<common::BitIntlB>();
-  obj->regKernel<common::BitDeintlB>();
-  obj->regKernel<cheetah::CommonTypeB>();
-  obj->regKernel<cheetah::CastTypeB>();
-  obj->regKernel<cheetah::ZeroB>();
-  obj->regKernel<cheetah::B2P>();
-  obj->regKernel<cheetah::P2B>();
-  obj->regKernel<cheetah::A2B>();
-  obj->regKernel<cheetah::B2A>();
-
-  obj->regKernel<cheetah::AndBP>();
-  obj->regKernel<cheetah::AndBB>();
-  obj->regKernel<cheetah::XorBP>();
-  obj->regKernel<cheetah::XorBB>();
-  obj->regKernel<cheetah::LShiftB>();
-  obj->regKernel<cheetah::RShiftB>();
-  obj->regKernel<cheetah::ARShiftB>();
-  obj->regKernel<cheetah::BitrevB>();
-  obj->regKernel<cheetah::RandA>();
-
-  return obj;
+  ctx->prot()->regKernel<cheetah::P2A>();
+  ctx->prot()->regKernel<cheetah::A2P>();
+  ctx->prot()->regKernel<cheetah::NotA>();
+  ctx->prot()->regKernel<cheetah::AddAP>();
+  ctx->prot()->regKernel<cheetah::AddAA>();
+  ctx->prot()->regKernel<cheetah::MulAP>();
+  ctx->prot()->regKernel<cheetah::MulAA>();
+  ctx->prot()->regKernel<cheetah::MulA1B>();
+  ctx->prot()->regKernel<cheetah::EqualAA>();
+  ctx->prot()->regKernel<cheetah::EqualAP>();
+  ctx->prot()->regKernel<cheetah::MatMulAP>();
+  ctx->prot()->regKernel<cheetah::MatMulAA>();
+  ctx->prot()->regKernel<cheetah::Conv2DAA>();
+  ctx->prot()->regKernel<cheetah::LShiftA>();
+  ctx->prot()->regKernel<cheetah::TruncA>();
+  ctx->prot()->regKernel<cheetah::TruncAWithSign>();
+  ctx->prot()->regKernel<cheetah::MsbA2B>();
+
+  ctx->prot()->regKernel<cheetah::CommonTypeB>();
+  ctx->prot()->regKernel<cheetah::CastTypeB>();
+  ctx->prot()->regKernel<cheetah::B2P>();
+  ctx->prot()->regKernel<cheetah::P2B>();
+  ctx->prot()->regKernel<cheetah::A2B>();
+  ctx->prot()->regKernel<cheetah::B2A>();
+
+  ctx->prot()->regKernel<cheetah::AndBP>();
+  ctx->prot()->regKernel<cheetah::AndBB>();
+  ctx->prot()->regKernel<cheetah::XorBP>();
+  ctx->prot()->regKernel<cheetah::XorBB>();
+  ctx->prot()->regKernel<cheetah::LShiftB>();
+  ctx->prot()->regKernel<cheetah::RShiftB>();
+  ctx->prot()->regKernel<cheetah::ARShiftB>();
+  ctx->prot()->regKernel<cheetah::BitrevB>();
+  ctx->prot()->regKernel<cheetah::RandA>();
+}
+
+std::unique_ptr<SPUContext> makeCheetahProtocol(
+    const RuntimeConfig& conf,
+    const std::shared_ptr<yacl::link::Context>& lctx) {
+  semi2k::registerTypes();
+
+  auto ctx = std::make_unique<SPUContext>(conf, lctx);
+
+  regCheetahProtocol(ctx.get(), lctx);
+
+  return ctx;
 }
 
 }  // namespace spu::mpc
diff --git a/libspu/mpc/cheetah/protocol.h b/libspu/mpc/cheetah/protocol.h
index 9a222ab2..10ea9166 100644
--- a/libspu/mpc/cheetah/protocol.h
+++ b/libspu/mpc/cheetah/protocol.h
@@ -16,11 +16,14 @@
 
 #include "yacl/link/link.h"
 
-#include "libspu/mpc/object.h"
+#include "libspu/core/context.h"
 
 namespace spu::mpc {
 
-std::unique_ptr<Object> makeCheetahProtocol(
+void regCheetahProtocol(SPUContext* ctx,
+                        const std::shared_ptr<yacl::link::Context>& lctx);
+
+std::unique_ptr<SPUContext> makeCheetahProtocol(
     const RuntimeConfig& conf,
     const std::shared_ptr<yacl::link::Context>& lctx);
 
diff --git a/libspu/mpc/cheetah/state.h b/libspu/mpc/cheetah/state.h
index 73834b63..ecfb3ed4 100644
--- a/libspu/mpc/cheetah/state.h
+++ b/libspu/mpc/cheetah/state.h
@@ -17,10 +17,10 @@
 #include <memory>
 #include <mutex>
 
+#include "libspu/core/object.h"
 #include "libspu/mpc/cheetah/arith/cheetah_dot.h"
 #include "libspu/mpc/cheetah/arith/cheetah_mul.h"
 #include "libspu/mpc/cheetah/ot/basic_ot_prot.h"
-#include "libspu/mpc/object.h"
 
 namespace spu::mpc::cheetah {
 
diff --git a/libspu/mpc/common/BUILD.bazel b/libspu/mpc/common/BUILD.bazel
index 79bfb74c..6b16c0d7 100644
--- a/libspu/mpc/common/BUILD.bazel
+++ b/libspu/mpc/common/BUILD.bazel
@@ -42,7 +42,7 @@ spu_cc_library(
     srcs = ["communicator.cc"],
     hdrs = ["communicator.h"],
     deps = [
-        "//libspu/mpc:object",
+        "//libspu/core:object",
         "//libspu/mpc/utils:ring_ops",
         "@yacl//yacl/link",
     ],
@@ -63,8 +63,8 @@ spu_cc_library(
     hdrs = ["prg_state.h"],
     deps = [
         "//libspu/core",
+        "//libspu/core:object",
         "//libspu/mpc:kernel",
-        "//libspu/mpc:object",
         "@yacl//yacl/crypto/tools:prg",
         "@yacl//yacl/crypto/utils:rand",
         "@yacl//yacl/link",
@@ -80,55 +80,6 @@ spu_cc_test(
     ],
 )
 
-spu_cc_library(
-    name = "ab_kernels",
-    srcs = ["ab_kernels.cc"],
-    hdrs = ["ab_kernels.h"],
-    deps = [
-        ":ab_api",
-        "//libspu/core:vectorize",
-        "//libspu/mpc:object",
-    ],
-)
-
-spu_cc_library(
-    name = "ab_api",
-    srcs = ["ab_api.cc"],
-    hdrs = ["ab_api.h"],
-    deps = [
-        ":pub2k",
-        "//libspu/core:vectorize",
-        "//libspu/mpc:api",
-        "//libspu/mpc:object",
-    ],
-)
-
-spu_cc_library(
-    name = "api_test_params",
-    hdrs = ["api_test_params.h"],
-    deps = [
-        "//libspu/mpc:object",
-        "@yacl//yacl/link",
-    ],
-)
-
-spu_cc_library(
-    name = "ab_api_test",
-    testonly = 1,
-    srcs = ["ab_api_test.cc"],
-    hdrs = ["ab_api_test.h"],
-    deps = [
-        ":ab_api",
-        ":api_test_params",
-        ":pub2k",
-        "//libspu/mpc:object",
-        "//libspu/mpc/common:communicator",
-        "//libspu/mpc/utils:simulate",
-        "@com_google_googletest//:gtest",
-    ],
-    alwayslink = True,
-)
-
 spu_cc_library(
     name = "prg_tensor",
     hdrs = ["prg_tensor.h"],
diff --git a/libspu/mpc/common/ab_api.cc b/libspu/mpc/common/ab_api.cc
deleted file mode 100644
index 13b04d41..00000000
--- a/libspu/mpc/common/ab_api.cc
+++ /dev/null
@@ -1,829 +0,0 @@
-// Copyright 2021 Ant Group Co., Ltd.
-//
-// 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 "libspu/mpc/common/ab_api.h"
-
-#include <future>
-
-#include "libspu/core/parallel_utils.h"
-#include "libspu/core/trace.h"
-#include "libspu/mpc/common/pub2k.h"
-
-namespace spu::mpc {
-namespace {
-
-// TODO(jint) may be we should move tiling to a `tiling` layer or dialect.
-ArrayRef block_par_unary(KernelEvalContext* ctx, std::string_view fn_name,
-                         const ArrayRef& in) {
-  const int64_t kBlockSize = kMinTaskSize;
-  if (!ctx->caller()->hasLowCostFork() || in.numel() <= kBlockSize) {
-    return ctx->caller()->call(fn_name, in);
-  }
-
-  std::string kBindName(fn_name);
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
-  auto* obj = ctx->caller();
-  std::vector<std::unique_ptr<Object>> sub_objs;
-
-  const int64_t numBlocks =
-      in.numel() / kBlockSize + ((in.numel() % kBlockSize) != 0 ? 1 : 0);
-
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    sub_objs.push_back(obj->fork());
-  }
-
-  std::vector<std::future<ArrayRef>> futures;
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    futures.push_back(std::async(
-        [&](int64_t index) {
-          int64_t begin = index * kBlockSize;
-          int64_t end = std::min(begin + kBlockSize, in.numel());
-
-          return sub_objs[index]->call(fn_name, in.slice(begin, end));
-        },
-        blk_idx));
-  }
-
-  std::vector<ArrayRef> out_slices;
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    out_slices.push_back(futures[blk_idx].get());
-  }
-
-  // Assume out.numel = in.numel
-  ArrayRef out(out_slices[0].eltype(), in.numel());
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    int64_t begin = blk_idx * kBlockSize;
-    int64_t end = std::min(begin + kBlockSize, in.numel());
-    std::memcpy(&out.at(begin), &out_slices[blk_idx].at(0),
-                (end - begin) * out.elsize());
-  }
-
-  return out;
-}
-
-ArrayRef block_par_unary_with_size(KernelEvalContext* ctx,
-                                   std::string_view fn_name, const ArrayRef& in,
-                                   size_t bits) {
-  const int64_t kBlockSize = kMinTaskSize;
-  if (!ctx->caller()->hasLowCostFork() || in.numel() <= kBlockSize) {
-    return ctx->caller()->call(fn_name, in, bits);
-  }
-
-  std::string kBindName(fn_name);
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
-  auto* obj = ctx->caller();
-  std::vector<std::unique_ptr<Object>> sub_objs;
-
-  const int64_t numBlocks =
-      in.numel() / kBlockSize + ((in.numel() % kBlockSize) != 0 ? 1 : 0);
-
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    sub_objs.push_back(obj->fork());
-  }
-
-  std::vector<std::future<ArrayRef>> futures;
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    futures.push_back(std::async(
-        [&](int64_t index) {
-          int64_t begin = index * kBlockSize;
-          int64_t end = std::min(begin + kBlockSize, in.numel());
-          return sub_objs[index]->call(fn_name, in.slice(begin, end), bits);
-        },
-        blk_idx));
-  }
-
-  std::vector<ArrayRef> out_slices;
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    out_slices.push_back(futures[blk_idx].get());
-  }
-
-  // Assume out.numel = in.numel
-  ArrayRef out(out_slices[0].eltype(), in.numel());
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    int64_t begin = blk_idx * kBlockSize;
-    int64_t end = std::min(begin + kBlockSize, in.numel());
-    std::memcpy(&out.at(begin), &out_slices[blk_idx].at(0),
-                (end - begin) * out.elsize());
-  }
-
-  return out;
-}
-
-ArrayRef block_par_binary(KernelEvalContext* ctx, std::string_view fn_name,
-                          const ArrayRef& lhs, const ArrayRef& rhs) {
-  const int64_t kBlockSize = kMinTaskSize;
-  SPU_ENFORCE(lhs.numel() == rhs.numel());
-  if (!ctx->caller()->hasLowCostFork() || lhs.numel() <= kBlockSize) {
-    return ctx->caller()->call(fn_name, lhs, rhs);
-  }
-
-  const int64_t numel = lhs.numel();
-
-  std::string kBindName(fn_name);
-  SPU_TRACE_MPC_LEAF(ctx, lhs);
-
-  auto* obj = ctx->caller();
-  std::vector<std::unique_ptr<Object>> sub_objs;
-
-  const int64_t numBlocks =
-      numel / kBlockSize + ((numel % kBlockSize) != 0 ? 1 : 0);
-
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    sub_objs.push_back(obj->fork());
-  }
-
-  std::vector<std::future<ArrayRef>> futures;
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    futures.push_back(std::async(
-        [&](int64_t index) {
-          int64_t begin = index * kBlockSize;
-          int64_t end = std::min(begin + kBlockSize, numel);
-
-          return sub_objs[index]->call(fn_name, lhs.slice(begin, end),
-                                       rhs.slice(begin, end));
-        },
-        blk_idx));
-  }
-
-  std::vector<ArrayRef> out_slices;
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    out_slices.push_back(futures[blk_idx].get());
-  }
-
-  // Assume out.numel = numel
-  ArrayRef out(out_slices[0].eltype(), numel);
-  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
-    int64_t begin = blk_idx * kBlockSize;
-    int64_t end = std::min(begin + kBlockSize, numel);
-    std::memcpy(&out.at(begin), &out_slices[blk_idx].at(0),
-                (end - begin) * out.elsize());
-  }
-
-  return out;
-}
-
-ArrayRef _Lazy2B(KernelEvalContext* ctx, const ArrayRef& in) {
-  if (in.eltype().isa<AShare>()) {
-    return block_par_unary(ctx, "a2b", in);
-  } else {
-    SPU_ENFORCE(in.eltype().isa<BShare>());
-    return in;
-  }
-}
-
-ArrayRef _Lazy2A(KernelEvalContext* ctx, const ArrayRef& in) {
-  if (in.eltype().isa<BShare>()) {
-    return block_par_unary(ctx, "b2a", in);
-  } else {
-    SPU_ENFORCE(in.eltype().isa<AShare>(), "expect AShare, got {}",
-                in.eltype());
-    return in;
-  }
-}
-
-// NOLINTBEGIN(bugprone-reserved-identifier)
-#define _LAZY_AB ctx->getState<ABProtState>()->lazy_ab
-
-#define _2A(x) _Lazy2A(ctx, x)
-#define _2B(x) _Lazy2B(ctx, x)
-
-#define _IsA(x) x.eltype().isa<AShare>()
-#define _IsB(x) x.eltype().isa<BShare>()
-#define _IsP(x) x.eltype().isa<Public>()
-#define _NBits(x) x.eltype().as<BShare>()->nbits()
-
-#define _A2P(x) ctx->caller()->call("a2p", x)
-#define _P2A(x) ctx->caller()->call("p2a", x)
-#define _NotA(x) ctx->caller()->call("not_a", x)
-#define _AddAP(lhs, rhs) ctx->caller()->call("add_ap", lhs, rhs)
-#define _AddAA(lhs, rhs) ctx->caller()->call("add_aa", lhs, rhs)
-#define _MulAP(lhs, rhs) ctx->caller()->call("mul_ap", lhs, rhs)
-#define _MulAA(lhs, rhs) block_par_binary(ctx, "mul_aa", lhs, rhs)
-#define _MulA1B(lhs, rhs) block_par_binary(ctx, "mul_a1b", lhs, rhs)
-#define _LShiftA(in, bits) ctx->caller()->call("lshift_a", in, bits)
-#define _TruncA(in, bits) block_par_unary_with_size(ctx, "trunc_a", in, bits)
-#define _MatMulAP(A, B, M, N, K) ctx->caller()->call("mmul_ap", A, B, M, N, K)
-#define _MatMulAA(A, B, M, N, K) ctx->caller()->call("mmul_aa", A, B, M, N, K)
-#define _B2P(x) ctx->caller()->call("b2p", x)
-#define _P2B(x) ctx->caller()->call("p2b", x)
-#define _A2B(x) block_par_unary(ctx, "a2b", x)
-#define _B2A(x) block_par_unary(ctx, "b2a", x)
-#define _NotB(x) ctx->caller()->call("not_b", x)
-#define _AndBP(lhs, rhs) ctx->caller()->call("and_bp", lhs, rhs)
-#define _AndBB(lhs, rhs) block_par_binary(ctx, "and_bb", lhs, rhs)
-#define _XorBP(lhs, rhs) ctx->caller()->call("xor_bp", lhs, rhs)
-#define _XorBB(lhs, rhs) ctx->caller()->call("xor_bb", lhs, rhs)
-#define _LShiftB(in, bits) ctx->caller()->call("lshift_b", in, bits)
-#define _RShiftB(in, bits) ctx->caller()->call("rshift_b", in, bits)
-#define _ARShiftB(in, bits) ctx->caller()->call("arshift_b", in, bits)
-#define _BitrevB(in, start, end) ctx->caller()->call("bitrev_b", in, start, end)
-#define _MsbA(in) block_par_unary(ctx, "msb_a2b", in)
-#define _RandA(size) ctx->caller()->call("rand_a", size)
-#define _RandB(size) ctx->caller()->call("rand_b", size)
-#define _EqualAP(lhs, rhs) block_par_binary(ctx, "equal_ap", lhs, rhs)
-#define _EqualAA(lhs, rhs) block_par_binary(ctx, "equal_aa", lhs, rhs)
-
-// NOLINTEND(bugprone-reserved-identifier)
-
-class ABProtState : public State {
- public:
-  static constexpr char kBindName[] = "ABProtState";
-
-  bool lazy_ab = true;
-
-  ABProtState() = default;
-  explicit ABProtState(bool lazy) : lazy_ab(lazy) {}
-
-  bool hasLowCostFork() const override { return true; }
-
-  std::unique_ptr<State> fork() override {
-    return std::make_unique<ABProtState>(lazy_ab);
-  }
-};
-
-class ABProtCommonTypeS : public Kernel {
- public:
-  static constexpr char kBindName[] = "common_type_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  void evaluate(KernelEvalContext* ctx) const override {
-    const Type& lhs = ctx->getParam<Type>(0);
-    const Type& rhs = ctx->getParam<Type>(1);
-
-    SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-
-    if (lhs.isa<AShare>() && rhs.isa<AShare>()) {
-      SPU_ENFORCE(lhs == rhs, "expect same, got lhs={}, rhs={}", lhs, rhs);
-      ctx->setOutput(lhs);
-    } else if (lhs.isa<AShare>() && rhs.isa<BShare>()) {
-      ctx->setOutput(lhs);
-    } else if (lhs.isa<BShare>() && rhs.isa<AShare>()) {
-      ctx->setOutput(rhs);
-    } else if (lhs.isa<BShare>() && rhs.isa<BShare>()) {
-      ctx->setOutput(common_type_b(ctx->caller(), lhs, rhs));
-    } else {
-      SPU_THROW("should not be here, lhs={}, rhs={}", lhs, rhs);
-    }
-  }
-};
-
-class ABProtCastTypeS : public Kernel {
- public:
-  static constexpr char kBindName[] = "cast_type_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  void evaluate(KernelEvalContext* ctx) const override {
-    const auto& frm = ctx->getParam<ArrayRef>(0);
-    const auto& to_type = ctx->getParam<Type>(1);
-
-    SPU_TRACE_MPC_DISP(ctx, frm, to_type);
-
-    if (frm.eltype().isa<AShare>() && to_type.isa<AShare>()) {
-      SPU_ENFORCE(frm.eltype() == to_type,
-                  "expect same, got frm={}, to_type={}", frm, to_type);
-      // do nothing.
-      ctx->setOutput(frm);
-    } else if (frm.eltype().isa<AShare>() && to_type.isa<BShare>()) {
-      ctx->setOutput(_A2B(frm));
-    } else if (frm.eltype().isa<BShare>() && to_type.isa<AShare>()) {
-      ctx->setOutput(_B2A(frm));
-    } else if (frm.eltype().isa<BShare>() && to_type.isa<BShare>()) {
-      ctx->setOutput(cast_type_b(ctx->caller(), frm, to_type));
-    } else {
-      SPU_THROW("should not be here, frm={}, to_type={}", frm, to_type);
-    }
-  }
-};
-
-class ABProtP2S : public UnaryKernel {
- public:
-  static constexpr char kBindName[] = "p2s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in) const override {
-    SPU_TRACE_MPC_DISP(ctx, in);
-    return _P2A(in);
-  }
-};
-
-class ABProtS2P : public UnaryKernel {
- public:
-  static constexpr char kBindName[] = "s2p";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in) const override {
-    SPU_TRACE_MPC_DISP(ctx, in);
-    if (_IsA(in)) {
-      return _A2P(in);
-    } else {
-      SPU_ENFORCE(_IsB(in));
-      return _B2P(in);
-    }
-  }
-};
-
-class ABProtRandS : public Kernel {
- public:
-  static constexpr char kBindName[] = "rand_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  void evaluate(KernelEvalContext* ctx) const override {
-    const size_t size = ctx->getParam<size_t>(0);
-
-    // ArrayRef proc(KernelEvalContext* ctx, size_t size) const override {
-    SPU_TRACE_MPC_DISP(ctx, size);
-
-    // always return random a share
-    ctx->setOutput(_RandA(size));
-  }
-};
-
-class ABProtNotS : public UnaryKernel {
- public:
-  static constexpr char kBindName[] = "not_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in) const override {
-    SPU_TRACE_MPC_DISP(ctx, in);
-    if (_LAZY_AB) {
-      // TODO: Both A&B could handle not(invert).
-      // if (in.eltype().isa<BShare>()) {
-      //  return _NotB(in);
-      //} else {
-      //  SPU_ENFORCE(in.eltype().isa<AShare>());
-      //  return _NotA(in);
-      //}
-      return _NotA(_2A(in));
-    }
-    return _NotA(in);
-  }
-};
-
-class ABProtAddSP : public BinaryKernel {
- public:
-  static constexpr char kBindName[] = "add_sp";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-    if (_LAZY_AB) {
-      return _AddAP(_2A(lhs), rhs);
-    }
-    return _AddAP(lhs, rhs);
-  }
-};
-
-class ABProtAddSS : public BinaryKernel {
- public:
-  static constexpr char kBindName[] = "add_ss";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-    if (_LAZY_AB) {
-      return _AddAA(_2A(lhs), _2A(rhs));
-    }
-    return _AddAA(lhs, rhs);
-  }
-};
-
-class ABProtMulSP : public BinaryKernel {
- public:
-  static constexpr char kBindName[] = "mul_sp";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-    if (_LAZY_AB) {
-      return _MulAP(_2A(lhs), rhs);
-    }
-    return _MulAP(lhs, rhs);
-  }
-};
-
-class ABProtMulSS : public BinaryKernel {
-  static bool hasMulA1B(KernelEvalContext* ctx) {
-    return ctx->hasKernel("mul_a1b");
-  }
-
- public:
-  static constexpr char kBindName[] = "mul_ss";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-
-    if (hasMulA1B(ctx) && _IsA(rhs) && _IsB(lhs) && _NBits(lhs) == 1) {
-      return _MulA1B(rhs, lhs);
-    }
-    if (hasMulA1B(ctx) && _IsA(lhs) && _IsB(rhs) && _NBits(rhs) == 1) {
-      return _MulA1B(lhs, rhs);
-    }
-    if (_LAZY_AB) {
-      // NOTE(juhou): Multiplication of two bits
-      if (_IsB(lhs) && _NBits(lhs) == 1 && _IsB(rhs) && _NBits(rhs) == 1) {
-        return _AndBB(lhs, rhs);
-      }
-      return _MulAA(_2A(lhs), _2A(rhs));
-    }
-    return _MulAA(lhs, rhs);
-  }
-};
-
-class ABProtMatMulSP : public MatmulKernel {
- public:
-  static constexpr char kBindName[] = "mmul_sp";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& a, const ArrayRef& b,
-                size_t m, size_t n, size_t k) const override {
-    SPU_TRACE_MPC_DISP(ctx, a, b);
-    if (_LAZY_AB) {
-      return _MatMulAP(_2A(a), b, m, n, k);
-    }
-    return _MatMulAP(a, b, m, n, k);
-  }
-};
-
-class ABProtMatMulSS : public MatmulKernel {
- public:
-  static constexpr char kBindName[] = "mmul_ss";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& a, const ArrayRef& b,
-                size_t m, size_t n, size_t k) const override {
-    SPU_TRACE_MPC_DISP(ctx, a, b);
-    if (_LAZY_AB) {
-      return _MatMulAA(_2A(a), _2A(b), m, n, k);
-    }
-    return _MatMulAA(a, b, m, n, k);
-  }
-};
-
-class ABProtAndSP : public BinaryKernel {
- public:
-  static constexpr char kBindName[] = "and_sp";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-    if (_LAZY_AB) {
-      return _AndBP(_2B(lhs), rhs);
-    }
-    return _B2A(_AndBP(_A2B(lhs), rhs));
-  }
-};
-
-class ABProtAndSS : public BinaryKernel {
- public:
-  static constexpr char kBindName[] = "and_ss";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-    if (_LAZY_AB) {
-      return _AndBB(_2B(lhs), _2B(rhs));
-    }
-    return _B2A(_AndBB(_A2B(lhs), _A2B(rhs)));
-  }
-};
-
-class ABProtXorSP : public BinaryKernel {
- public:
-  static constexpr char kBindName[] = "xor_sp";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-    if (_LAZY_AB) {
-      return _XorBP(_2B(lhs), rhs);
-    }
-    return _B2A(_XorBP(_A2B(lhs), rhs));
-  }
-};
-
-class ABProtXorSS : public BinaryKernel {
- public:
-  static constexpr char kBindName[] = "xor_ss";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-    if (_LAZY_AB) {
-      return _XorBB(_2B(lhs), _2B(rhs));
-    }
-    return _B2A(_XorBB(_A2B(lhs), _A2B(rhs)));
-  }
-};
-
-class ABProtEqualSS : public Kernel {
- public:
-  static constexpr char kBindName[] = "equal_ss";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  void evaluate(KernelEvalContext* ctx) const override {
-    const auto& x = ctx->getParam<ArrayRef>(0);
-    const auto& y = ctx->getParam<ArrayRef>(1);
-
-    SPU_TRACE_MPC_DISP(ctx, x, y);
-
-    // TODO: use mandatory lazyAB pattern.
-    SPU_ENFORCE(_LAZY_AB);
-
-    ctx->setOutput(std::nullopt);
-
-    // try fast path
-    // TODO: use cost model instead of hand-coded priority.
-    if (_IsA(x) && _IsA(y)) {
-      if (ctx->hasKernel("equal_aa")) {
-        ctx->setOutput(ctx->call("equal_aa", x, y));
-      }
-    } else if (_IsB(x) && _IsB(y)) {
-      if (ctx->hasKernel("equal_bb")) {
-        ctx->setOutput(ctx->call("equal_bb", x, y));
-      }
-    } else {
-      // mixed A and B
-      SPU_ENFORCE((_IsA(x) && _IsB(y)) || (_IsB(x) && _IsA(y)));
-
-      if (ctx->hasKernel("equal_aa")) {
-        ctx->setOutput(ctx->call("equal_aa", _2A(x), _2A(y)));
-      }
-
-      if (ctx->hasKernel("equal_bb")) {
-        ctx->setOutput(ctx->call("equal_bb", _2B(x), _2B(y)));
-      }
-    }
-
-    // fall through, no kernel found.
-  }
-};
-
-class ABProtEqualSP : public Kernel {
- public:
-  static constexpr char kBindName[] = "equal_sp";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  void evaluate(KernelEvalContext* ctx) const override {
-    const auto& x = ctx->getParam<ArrayRef>(0);
-    const auto& y = ctx->getParam<ArrayRef>(1);
-
-    SPU_TRACE_MPC_DISP(ctx, x, y);
-
-    // TODO: use mandatory lazyAB.
-    SPU_ENFORCE(_LAZY_AB);
-
-    ctx->setOutput(std::nullopt);
-
-    if (_IsA(x) && ctx->hasKernel("equal_ap")) {
-      ctx->setOutput(ctx->call("equal_ap", x, y));
-    } else if (_IsB(x) && ctx->hasKernel("equal_bp")) {
-      ctx->setOutput(ctx->call("equal_bp", x, y));
-    }
-  }
-};
-
-class ABProtLShiftS : public ShiftKernel {
- public:
-  static constexpr char kBindName[] = "lshift_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
-                size_t bits) const override {
-    SPU_TRACE_MPC_DISP(ctx, in, bits);
-    if (in.eltype().isa<AShare>()) {
-      return _LShiftA(in, bits);
-    } else if (in.eltype().isa<BShare>()) {
-      if (_LAZY_AB) {
-        return _LShiftB(in, bits);
-      } else {
-        return _B2A(_LShiftB(in, bits));
-      }
-    } else {
-      SPU_THROW("Unsupported type {}", in.eltype());
-    }
-  }
-};
-
-class ABProtRShiftS : public ShiftKernel {
- public:
-  static constexpr char kBindName[] = "rshift_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
-                size_t bits) const override {
-    SPU_TRACE_MPC_DISP(ctx, in, bits);
-    if (_LAZY_AB) {
-      return _RShiftB(_2B(in), bits);
-    }
-    return _B2A(_RShiftB(_A2B(in), bits));
-  }
-};
-
-class ABProtARShiftS : public ShiftKernel {
- public:
-  static constexpr char kBindName[] = "arshift_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
-                size_t bits) const override {
-    SPU_TRACE_MPC_DISP(ctx, in, bits);
-    if (_LAZY_AB) {
-      return _ARShiftB(_2B(in), bits);
-    }
-    return _B2A(_ARShiftB(_A2B(in), bits));
-  }
-};
-
-class ABProtTruncS : public ShiftKernel {
- public:
-  static constexpr char kBindName[] = "trunc_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
-                size_t bits) const override {
-    SPU_TRACE_MPC_DISP(ctx, in, bits);
-    if (_LAZY_AB) {
-      return _TruncA(_2A(in), bits);
-    }
-    return _TruncA(in, bits);
-  }
-};
-
-class ABProtBitrevS : public Kernel {
- public:
-  static constexpr char kBindName[] = "bitrev_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<ArrayRef>(0),
-                        ctx->getParam<size_t>(1), ctx->getParam<size_t>(2)));
-  }
-  static ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in, size_t start,
-                       size_t end) {
-    SPU_TRACE_MPC_DISP(ctx, in, start, end);
-    if (_LAZY_AB) {
-      return _BitrevB(_2B(in), start, end);
-    }
-    return _B2A(_BitrevB(_A2B(in), start, end));
-  }
-};
-
-class ABProtMsbS : public UnaryKernel {
- public:
-  static constexpr char kBindName[] = "msb_s";
-
-  Kind kind() const override { return Kind::Dynamic; }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in) const override {
-    SPU_TRACE_MPC_DISP(ctx, in);
-    const auto field = in.eltype().as<Ring2k>()->field();
-    if (ctx->hasKernel("msb_a2b")) {
-      if (_LAZY_AB) {
-        if (in.eltype().isa<BShare>()) {
-          return _RShiftB(in, SizeOf(field) * 8 - 1);
-        } else {
-          // fast path, directly apply msb in AShare, result a BShare.
-          return _MsbA(in);
-        }
-      } else {
-        // Do it in AShare domain, and convert back to AShare.
-        return _B2A(_MsbA(in));
-      }
-    } else {
-      if (_LAZY_AB) {
-        return _RShiftB(_2B(in), SizeOf(field) * 8 - 1);
-      }
-      return _B2A(_RShiftB(_A2B(in), SizeOf(field) * 8 - 1));
-    }
-  }
-};
-
-}  // namespace
-
-Type common_type_b(Object* ctx, const Type& a, const Type& b) {
-  return ctx->call<Type>("common_type_b", a, b);
-}
-
-ArrayRef cast_type_b(Object* ctx, const ArrayRef& a, const Type& to_type) {
-  return ctx->call("cast_type_b", a, to_type);
-}
-
-ArrayRef zero_a(Object* ctx, size_t sz) { return ctx->call("zero_a", sz); }
-
-ArrayRef rand_a(Object* ctx, size_t sz) { return ctx->call("rand_a", sz); }
-
-ArrayRef zero_b(Object* ctx, size_t sz) { return ctx->call("zero_b", sz); }
-
-ArrayRef rand_b(Object* ctx, size_t sz) { return ctx->call("rand_b", sz); }
-
-SPU_MPC_DEF_UNARY_OP(a2p)
-SPU_MPC_DEF_UNARY_OP(p2a)
-SPU_MPC_DEF_UNARY_OP(msb_a2b)
-SPU_MPC_DEF_UNARY_OP(not_a)
-SPU_MPC_DEF_BINARY_OP(add_ap)
-SPU_MPC_DEF_BINARY_OP(add_aa)
-SPU_MPC_DEF_BINARY_OP(mul_ap)
-SPU_MPC_DEF_BINARY_OP(mul_aa)
-SPU_MPC_DEF_BINARY_OP(mul_a1b)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(lshift_a)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(trunc_a)
-SPU_MPC_DEF_MMUL(mmul_ap)
-SPU_MPC_DEF_MMUL(mmul_aa)
-
-SPU_MPC_DEF_UNARY_OP(b2p)
-SPU_MPC_DEF_UNARY_OP(p2b)
-SPU_MPC_DEF_UNARY_OP(a2b)
-SPU_MPC_DEF_UNARY_OP(b2a)
-SPU_MPC_DEF_BINARY_OP(and_bp)
-SPU_MPC_DEF_BINARY_OP(and_bb)
-SPU_MPC_DEF_BINARY_OP(xor_bp)
-SPU_MPC_DEF_BINARY_OP(xor_bb)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(lshift_b)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(rshift_b)
-SPU_MPC_DEF_UNARY_OP_WITH_SIZE(arshift_b)
-SPU_MPC_DEF_UNARY_OP_WITH_2SIZE(bitrev_b);
-SPU_MPC_DEF_BINARY_OP(add_bb)
-
-ArrayRef bitintl_b(Object* ctx, const ArrayRef& in, size_t stride) {
-  return ctx->call("bitintl_b", in, stride);
-}
-
-ArrayRef bitdeintl_b(Object* ctx, const ArrayRef& in, size_t stride) {
-  return ctx->call("bitdeintl_b", in, stride);
-}
-
-void regABKernels(Object* obj) {
-  obj->addState<ABProtState>();
-
-  obj->regKernel<ABProtCommonTypeS>();
-  obj->regKernel<ABProtCastTypeS>();
-  obj->regKernel<ABProtP2S>();
-  obj->regKernel<ABProtS2P>();
-  obj->regKernel<ABProtRandS>();
-  obj->regKernel<ABProtNotS>();
-  obj->regKernel<ABProtAddSP>();
-  obj->regKernel<ABProtAddSS>();
-  obj->regKernel<ABProtMulSP>();
-  obj->regKernel<ABProtMulSS>();
-  obj->regKernel<ABProtMatMulSP>();
-  obj->regKernel<ABProtMatMulSS>();
-  obj->regKernel<ABProtAndSP>();
-  obj->regKernel<ABProtAndSS>();
-  obj->regKernel<ABProtXorSP>();
-  obj->regKernel<ABProtXorSS>();
-  obj->regKernel<ABProtEqualSS>();
-  obj->regKernel<ABProtEqualSP>();
-  obj->regKernel<ABProtLShiftS>();
-  obj->regKernel<ABProtRShiftS>();
-  obj->regKernel<ABProtARShiftS>();
-  obj->regKernel<ABProtTruncS>();
-  obj->regKernel<ABProtBitrevS>();
-  obj->regKernel<ABProtMsbS>();
-}
-
-}  // namespace spu::mpc
diff --git a/libspu/mpc/common/ab_api.h b/libspu/mpc/common/ab_api.h
deleted file mode 100644
index a5875ab6..00000000
--- a/libspu/mpc/common/ab_api.h
+++ /dev/null
@@ -1,108 +0,0 @@
-// Copyright 2021 Ant Group Co., Ltd.
-//
-// 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.
-
-#pragma once
-
-#include "libspu/mpc/api.h"
-#include "libspu/mpc/kernel.h"
-#include "libspu/mpc/object.h"
-
-namespace spu::mpc {
-
-ArrayRef a2p(Object* ctx, const ArrayRef&);
-ArrayRef p2a(Object* ctx, const ArrayRef&);
-ArrayRef msb_a2b(Object* ctx, const ArrayRef&);
-
-ArrayRef zero_a(Object* ctx, size_t);
-ArrayRef rand_a(Object* ctx, size_t);
-ArrayRef rand_b(Object* ctx, size_t);
-
-ArrayRef not_a(Object* ctx, const ArrayRef&);
-
-ArrayRef add_ap(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef add_aa(Object* ctx, const ArrayRef&, const ArrayRef&);
-
-ArrayRef mul_ap(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef mul_aa(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef mul_a1b(Object* ctx, const ArrayRef&, const ArrayRef&);
-
-ArrayRef lshift_a(Object* ctx, const ArrayRef&, size_t);
-ArrayRef trunc_a(Object* ctx, const ArrayRef&, size_t);
-
-ArrayRef mmul_ap(Object* ctx, const ArrayRef&, const ArrayRef&, size_t, size_t,
-                 size_t);
-ArrayRef mmul_aa(Object* ctx, const ArrayRef&, const ArrayRef&, size_t, size_t,
-                 size_t);
-
-ArrayRef zero_b(Object* ctx, size_t);
-
-Type common_type_b(Object* ctx, const Type& a, const Type& b);
-ArrayRef cast_type_b(Object* ctx, const ArrayRef& a, const Type& to_type);
-
-ArrayRef b2p(Object* ctx, const ArrayRef&);
-ArrayRef p2b(Object* ctx, const ArrayRef&);
-
-ArrayRef a2b(Object* ctx, const ArrayRef&);
-ArrayRef b2a(Object* ctx, const ArrayRef&);
-
-ArrayRef and_bp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef and_bb(Object* ctx, const ArrayRef&, const ArrayRef&);
-
-ArrayRef xor_bp(Object* ctx, const ArrayRef&, const ArrayRef&);
-ArrayRef xor_bb(Object* ctx, const ArrayRef&, const ArrayRef&);
-
-ArrayRef lshift_b(Object* ctx, const ArrayRef&, size_t);
-ArrayRef rshift_b(Object* ctx, const ArrayRef&, size_t);
-ArrayRef arshift_b(Object* ctx, const ArrayRef&, size_t);
-
-// Bit reverse for binary share.
-ArrayRef bitrev_b(Object* ctx, const ArrayRef&, size_t, size_t);
-
-// TODO: maybe we should add more general PDEP/PEXT instruction.
-
-/// Bit interleave for binary share.
-//
-// Interleave bits of input, so the upper bits of input are in the even
-// positions and lower bits in the odd. Also called Morton Number.
-//
-//   abcdXYZW -> aXbYcZdW
-//
-// stride represent the log shift of the interleaved distance.
-//
-//   00001111 -> 01010101     stride = 0
-//   00001111 -> 00110011     stride = 1
-//   00001111 -> 00001111     stride = 2
-//
-ArrayRef bitintl_b(Object* ctx, const ArrayRef& in, size_t stride);
-
-// Bit de-interleave for binary share.
-//
-// The reverse bit interleave method, put the even bits at lower half, and odd
-// bits at upper half.
-//
-//   aXbYcZdW -> abcdXYZW
-//
-// stride represent the log shift of the interleaved distance.
-//
-//   01010101 -> 00001111     stride = 0
-//   00110011 -> 00001111     stride = 1
-//   00001111 -> 00001111     stride = 2
-//
-ArrayRef bitdeintl_b(Object* ctx, const ArrayRef& in, size_t stride);
-
-ArrayRef add_bb(Object* ctx, const ArrayRef&, const ArrayRef&);
-
-void regABKernels(Object* obj);
-
-}  // namespace spu::mpc
diff --git a/libspu/mpc/common/ab_kernels.cc b/libspu/mpc/common/ab_kernels.cc
deleted file mode 100644
index 75618ad5..00000000
--- a/libspu/mpc/common/ab_kernels.cc
+++ /dev/null
@@ -1,306 +0,0 @@
-// Copyright 2022 Ant Group Co., Ltd.
-//
-// 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 "libspu/mpc/common/ab_kernels.h"
-
-#include "yacl/base/int128.h"
-
-#include "libspu/core/bit_utils.h"
-#include "libspu/core/trace.h"
-#include "libspu/core/vectorize.h"
-#include "libspu/mpc/api.h"
-#include "libspu/mpc/common/ab_api.h"
-
-namespace spu::mpc::common {
-namespace {
-
-constexpr std::array<uint128_t, 6> kBitIntlSwapMasks = {{
-    yacl::MakeUint128(0x2222222222222222, 0x2222222222222222),  // 4bit
-    yacl::MakeUint128(0x0C0C0C0C0C0C0C0C, 0x0C0C0C0C0C0C0C0C),  // 8bit
-    yacl::MakeUint128(0x00F000F000F000F0, 0x00F000F000F000F0),  // 16bit
-    yacl::MakeUint128(0x0000FF000000FF00, 0x0000FF000000FF00),  // 32bit
-    yacl::MakeUint128(0x00000000FFFF0000, 0x00000000FFFF0000),  // 64bit
-    yacl::MakeUint128(0x0000000000000000, 0xFFFFFFFF00000000),  // 128bit
-}};
-
-constexpr std::array<uint128_t, 6> kBitIntlKeepMasks = {{
-    yacl::MakeUint128(0x9999999999999999, 0x9999999999999999),  // 4bit
-    yacl::MakeUint128(0xC3C3C3C3C3C3C3C3, 0xC3C3C3C3C3C3C3C3),  // 8bit
-    yacl::MakeUint128(0xF00FF00FF00FF00F, 0xF00FF00FF00FF00F),  // 16bit
-    yacl::MakeUint128(0xFF0000FFFF0000FF, 0xFF0000FFFF0000FF),  // 32bit
-    yacl::MakeUint128(0xFFFF00000000FFFF, 0xFFFF00000000FFFF),  // 64bit
-    yacl::MakeUint128(0xFFFFFFFF00000000, 0x00000000FFFFFFFF),  // 128bit
-}};
-
-inline size_t numBits(const ArrayRef& in) {
-  return in.eltype().as<BShare>()->nbits();
-}
-
-inline ArrayRef setNumBits(const ArrayRef& in, size_t nbits) {
-  ArrayRef out = in;
-  out.eltype().as<BShare>()->setNbits(nbits);
-  return out;
-}
-
-}  // namespace
-
-// implement interleave by xor_bb/and_bp/shift_b, assume all these
-// instructions are communication free.
-void BitIntlB::evaluate(KernelEvalContext* ctx) const {
-  const auto& in = ctx->getParam<ArrayRef>(0);
-  const size_t stride = ctx->getParam<size_t>(1);
-  auto* obj = ctx->caller();
-
-  SPU_TRACE_MPC_LEAF(ctx, in, stride);
-
-  // algorithm:
-  //      0000000011111111
-  // swap     ^^^^^^^^
-  //      0000111100001111
-  // swap   ^^^^    ^^^^
-  //      0011001100110011
-  // swap  ^^  ^^  ^^  ^^
-  //      0101010101010101
-  const size_t nbits = in.eltype().as<BShare>()->nbits();
-  SPU_ENFORCE(absl::has_single_bit(nbits));
-
-  const size_t size = in.numel();
-  ArrayRef out = in.clone();
-  for (int64_t idx = Log2Ceil(nbits) - 2; idx >= static_cast<int64_t>(stride);
-       idx--) {
-    auto K = make_p(obj, kBitIntlKeepMasks[idx], size);
-    auto M = make_p(obj, kBitIntlSwapMasks[idx], size);
-    int64_t S = 1 << idx;
-    // out = (out & K) ^ ((out >> S) & M) ^ ((out & M) << S);
-    out = xor_bb(
-        obj,
-        xor_bb(obj, and_bp(obj, out, K), and_bp(obj, rshift_b(obj, out, S), M)),
-        lshift_b(obj, and_bp(obj, out, M), S));
-  }
-  out = setNumBits(out, numBits(in));
-  ctx->setOutput(out);
-}
-
-// implement interleave by xor_bb/and_bp/shift_b, assume all these
-// instructions are communication free.
-void BitDeintlB::evaluate(KernelEvalContext* ctx) const {
-  const auto& in = ctx->getParam<ArrayRef>(0);
-  const size_t stride = ctx->getParam<size_t>(1);
-  auto* obj = ctx->caller();
-
-  SPU_TRACE_MPC_LEAF(ctx, in, stride);
-
-  // algorithm:
-  //      0101010101010101
-  // swap  ^^  ^^  ^^  ^^
-  //      0011001100110011
-  // swap   ^^^^    ^^^^
-  //      0000111100001111
-  // swap     ^^^^^^^^
-  //      0000000011111111
-  const size_t nbits = in.eltype().as<BShare>()->nbits();
-  SPU_ENFORCE(absl::has_single_bit(nbits));
-  const size_t size = in.numel();
-
-  ArrayRef out = in.clone();
-  for (int64_t idx = stride; idx + 1 < Log2Ceil(nbits); idx++) {
-    auto K = make_p(obj, kBitIntlKeepMasks[idx], size);
-    auto M = make_p(obj, kBitIntlSwapMasks[idx], size);
-    int64_t S = 1 << idx;
-    // out = (out & K) ^ ((out >> S) & M) ^ ((out & M) << S);
-    out = xor_bb(
-        obj,
-        xor_bb(obj, and_bp(obj, out, K), and_bp(obj, rshift_b(obj, out, S), M)),
-        lshift_b(obj, and_bp(obj, out, M), S));
-  }
-  out = setNumBits(out, numBits(in));
-  ctx->setOutput(out);
-}
-
-namespace {
-
-// The kogge-stone adder.
-//
-// P stands for propagate, G stands for generate, where:
-//  (G0, P0) = (g0, p0)
-//  (Gi, Pi) = (gi, pi) o (Gi-1, Pi-1)
-//
-// The `o` here is:
-//  (G0, P0) o (G1, P1) = (G0 ^ (P0 & G1), P0 & P1)
-//
-// Latency log(k) + 1
-ArrayRef ppa_kogge_stone(Object* ctx, const ArrayRef& lhs, const ArrayRef& rhs,
-                         size_t nbits) {
-  // Generate p & g.
-  auto P = xor_bb(ctx, lhs, rhs);
-  auto G = and_bb(ctx, lhs, rhs);
-
-  for (int idx = 0; idx < Log2Ceil(nbits); ++idx) {
-    const size_t offset = 1UL << idx;
-    auto G1 = lshift_b(ctx, G, offset);
-    auto P1 = lshift_b(ctx, P, offset);
-
-    // P1 = P & P1
-    // G1 = G ^ (P & G1)
-    std::vector<ArrayRef> res = spu::vectorize(
-        {P, P}, {P1, G1}, [&](const ArrayRef& xx, const ArrayRef& yy) {
-          return and_bb(ctx, xx, yy);
-        });
-    P = std::move(res[0]);
-    G = xor_bb(ctx, G, res[1]);
-  }
-
-  // out = (G << 1) ^ p0
-  auto C = lshift_b(ctx, G, 1);
-  return xor_bb(ctx, xor_bb(ctx, lhs, rhs), C);
-}
-
-std::pair<ArrayRef, ArrayRef> bit_scatter(Object* ctx, const ArrayRef& in,
-                                          size_t stride) {
-  const size_t nbits = numBits(in);
-  SPU_ENFORCE(absl::has_single_bit(nbits), "unsupported {}", nbits);
-  auto out = bitdeintl_b(ctx, in, stride);
-
-  auto hi = rshift_b(ctx, out, nbits / 2);
-  auto mask =
-      make_p(ctx, (static_cast<uint128_t>(1) << (nbits / 2)) - 1, in.numel());
-  auto lo = and_bp(ctx, out, mask);
-  return std::make_pair(hi, lo);
-}
-
-ArrayRef bit_gather(Object* ctx, const ArrayRef& hi, const ArrayRef& lo,
-                    size_t stride) {
-  const size_t nbits = numBits(hi);
-  SPU_ENFORCE(absl::has_single_bit(nbits), "unsupported {}", nbits);
-  SPU_ENFORCE(nbits == numBits(lo), "nbits mismatch {}, {}", nbits,
-              numBits(lo));
-
-  auto out = xor_bb(ctx, lshift_b(ctx, hi, nbits), lo);
-  return bitintl_b(ctx, out, stride);
-}
-
-// The sklansky adder.
-ArrayRef ppa_sklansky(Object* ctx, ArrayRef const& lhs, ArrayRef const& rhs,
-                      size_t nbits) {
-  SPU_ENFORCE(lhs.numel() == rhs.numel());
-
-  constexpr std::array<uint128_t, 7> kSelMask = {{
-      yacl::MakeUint128(0xFFFFFFFFFFFFFFFF, 0xFFFFFFFFFFFFFFFF),  // invalid
-      yacl::MakeUint128(0xAAAAAAAAAAAAAAAA, 0xAAAAAAAAAAAAAAAA),  // 10101010
-      yacl::MakeUint128(0x8888888888888888, 0x8888888888888888),  // 10001000
-      yacl::MakeUint128(0x8080808080808080, 0x8080808080808080),  // 10000000
-      yacl::MakeUint128(0x8000800080008000, 0x8000800080008000),  // ...
-      yacl::MakeUint128(0x8000000080000000, 0x8000000080000000),  // ...
-      yacl::MakeUint128(0x8000000000000000, 0x8000000000000000),  // ...
-  }};
-
-  // Generate P & G.
-  auto P = xor_bb(ctx, lhs, rhs);
-  auto G = and_bb(ctx, lhs, rhs);
-
-  const size_t bit_width = numBits(lhs);
-  SPU_ENFORCE(bit_width == numBits(rhs), "nbits mismatch {}, {}", bit_width,
-              numBits(rhs));
-  for (int idx = 0; idx < Log2Ceil(nbits); ++idx) {
-    auto [Ph, Pl] = bit_scatter(ctx, P, idx);
-    auto [Gh, Gl] = bit_scatter(ctx, G, idx);
-    // SPU_ENFORCE(numBits(Ph) == bit_width / 2);
-    // SPU_ENFORCE(numBits(Pl) == bit_width / 2);
-    // SPU_ENFORCE(numBits(Gh) == bit_width / 2);
-    // SPU_ENFORCE(numBits(Gl) == bit_width / 2);
-
-    const auto s_mask = make_p(ctx, kSelMask[idx], lhs.numel());
-    auto Gs = and_bp(ctx, Gl, s_mask);
-    auto Ps = and_bp(ctx, Pl, s_mask);
-    for (int j = 0; j < idx; j++) {
-      Gs = xor_bb(ctx, Gs, rshift_b(ctx, Gs, 1 << j));
-      Ps = xor_bb(ctx, Ps, rshift_b(ctx, Ps, 1 << j));
-    }
-    // SPU_ENFORCE(numBits(Ps) == bit_width / 2);
-    // SPU_ENFORCE(numBits(Gs) == bit_width / 2);
-
-    // Ph = Ph & Ps
-    // Gh = Gh ^ (Ph & Gs)
-    std::vector<ArrayRef> PG = spu::vectorize(
-        {Ph, Ph}, {Ps, Gs}, [&](const ArrayRef& xx, const ArrayRef& yy) {
-          return and_bb(ctx, xx, yy);
-        });
-    Ph = std::move(PG[0]);
-    Gh = xor_bb(ctx, Gh, PG[1]);
-    // SPU_ENFORCE(numBits(Gh) == numBits(G) / 2);
-    // SPU_ENFORCE(numBits(Ph) == numBits(P) / 2);
-
-    P = bit_gather(ctx, Ph, Pl, idx);
-    G = bit_gather(ctx, Gh, Gl, idx);
-  }
-
-  // out = (G0 << 1) ^ p0
-  auto C = lshift_b(ctx, G, 1);
-  return xor_bb(ctx, xor_bb(ctx, lhs, rhs), C);
-}
-
-}  // namespace
-
-ArrayRef AddBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                     const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
-  const size_t nbits = numBits(lhs);
-  SPU_ENFORCE(nbits == numBits(rhs), "nbits mismatch {}!={}", nbits,
-              numBits(rhs));
-
-  switch (type_) {
-    case CircuitType::KoggeStone:
-      return ppa_kogge_stone(ctx->caller(), lhs, rhs, nbits);
-    case CircuitType::Sklansky:
-      return ppa_sklansky(ctx->caller(), lhs, rhs, nbits);
-    default:
-      SPU_THROW("unknown circuit type {}", static_cast<uint32_t>(type_));
-  }
-}
-
-// compute the k'th bit of x + y
-ArrayRef carry_out(Object* ctx, const ArrayRef& x, const ArrayRef& y,
-                   size_t k) {
-  // init P & G
-  auto P = xor_bb(ctx, x, y);
-  auto G = and_bb(ctx, x, y);
-
-  // Use kogge stone layout.
-  while (k > 1) {
-    if (k % 2 != 0) {
-      k += 1;
-      P = lshift_b(ctx, P, 1);
-      G = lshift_b(ctx, G, 1);
-    }
-    auto [P1, P0] = bit_scatter(ctx, P, 0);
-    auto [G1, G0] = bit_scatter(ctx, G, 0);
-
-    // Calculate next-level of P, G
-    //   P = P1 & P0
-    //   G = G1 | (P1 & G0)
-    //     = G1 ^ (P1 & G0)
-    std::vector<ArrayRef> v = vectorize(
-        {P0, G0}, {P1, P1}, [&](const ArrayRef& xx, const ArrayRef& yy) {
-          return and_bb(ctx, xx, yy);
-        });
-    P = std::move(v[0]);
-    G = xor_bb(ctx, G1, v[1]);
-    k >>= 1;
-  }
-
-  return G;
-}
-
-}  // namespace spu::mpc::common
diff --git a/libspu/mpc/common/ab_kernels.h b/libspu/mpc/common/ab_kernels.h
deleted file mode 100644
index b1665507..00000000
--- a/libspu/mpc/common/ab_kernels.h
+++ /dev/null
@@ -1,92 +0,0 @@
-// Copyright 2022 Ant Group Co., Ltd.
-//
-// 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.
-
-#pragma once
-
-#include "libspu/mpc/kernel.h"
-#include "libspu/mpc/utils/cexpr.h"
-
-namespace spu::mpc::common {
-
-class BitIntlB : public Kernel {
- public:
-  static constexpr char kBindName[] = "bitintl_b";
-
-  ce::CExpr latency() const override { return ce::Const(0); }
-
-  ce::CExpr comm() const override { return ce::Const(0); }
-
-  void evaluate(KernelEvalContext* ctx) const override;
-};
-
-class BitDeintlB : public Kernel {
- public:
-  static constexpr char kBindName[] = "bitdeintl_b";
-
-  ce::CExpr latency() const override { return ce::Const(0); }
-
-  ce::CExpr comm() const override { return ce::Const(0); }
-
-  void evaluate(KernelEvalContext* ctx) const override;
-};
-
-class AddBB : public BinaryKernel {
- public:
-  static constexpr char kBindName[] = "add_bb";
-
-  enum class CircuitType {
-    KoggeStone,
-    Sklansky,
-    Count,
-  };
-
- private:
-  CircuitType type_;
-
- public:
-  explicit AddBB(CircuitType type = CircuitType::KoggeStone) : type_(type) {}
-
-  ce::CExpr latency() const override {
-    // For both sklansky/ks
-    return Log(ce::K()) + ce::Const(1);
-  }
-
-  // TODO: depends on the and gate.
-  ce::CExpr comm() const override {
-    switch (type_) {
-      case CircuitType::KoggeStone: {
-        // FIXME: this is aby3 cost
-        // 1 * and        : k
-        // logk * 2 * and : logk * 2k
-        return Log(ce::K()) * ce::K() * 2 + ce::K();
-      }
-      case CircuitType::Sklansky: {
-        // FIXME: this is aby3 cost
-        // 1 * and    : k
-        // logk * and : logk * k
-        return Log(ce::K()) * ce::K() + ce::K();
-      }
-      default:
-        SPU_THROW("should not be here");
-    }
-  }
-
-  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
-                const ArrayRef& rhs) const override;
-};
-
-// compute the k'th bit of x + y
-ArrayRef carry_out(Object* ctx, const ArrayRef& x, const ArrayRef& y, size_t k);
-
-}  // namespace spu::mpc::common
diff --git a/libspu/mpc/common/communicator.h b/libspu/mpc/common/communicator.h
index 2a664338..d92ea12d 100644
--- a/libspu/mpc/common/communicator.h
+++ b/libspu/mpc/common/communicator.h
@@ -23,9 +23,9 @@
 #include "yacl/base/buffer.h"
 #include "yacl/link/link.h"
 
+#include "libspu/core/object.h"
 #include "libspu/core/parallel_utils.h"
 #include "libspu/core/prelude.h"
-#include "libspu/mpc/object.h"
 
 // This module defines the protocol comm pattern used for all
 // protocols.
diff --git a/libspu/mpc/common/prg_state.h b/libspu/mpc/common/prg_state.h
index e6c0d45c..fbeb9675 100644
--- a/libspu/mpc/common/prg_state.h
+++ b/libspu/mpc/common/prg_state.h
@@ -19,7 +19,7 @@
 #include "yacl/link/link.h"
 
 #include "libspu/core/array_ref.h"
-#include "libspu/mpc/object.h"
+#include "libspu/core/object.h"
 
 namespace spu::mpc {
 
diff --git a/libspu/mpc/common/pub2k.cc b/libspu/mpc/common/pub2k.cc
index 0349f520..e299a27d 100644
--- a/libspu/mpc/common/pub2k.cc
+++ b/libspu/mpc/common/pub2k.cc
@@ -34,26 +34,26 @@ class Pub2kMakeP : public Kernel {
 
   void evaluate(KernelEvalContext* ctx) const override {
     ctx->setOutput(
-        proc(ctx, ctx->getParam<uint128_t>(0), ctx->getParam<size_t>(1)));
+        proc(ctx, ctx->getParam<uint128_t>(0), ctx->getParam<Shape>(1)));
   }
 
-  static ArrayRef proc(KernelEvalContext* ctx, uint128_t init, size_t size) {
-    SPU_TRACE_MPC_LEAF(ctx, init);
+  static Value proc(KernelEvalContext* ctx, uint128_t init,
+                    const Shape& shape) {
     const auto field = ctx->getState<Z2kState>()->getDefaultField();
 
     const auto eltype = makeType<Pub2kTy>(field);
     auto buf = std::make_shared<yacl::Buffer>(1 * eltype.size());
-    ArrayRef res(buf,     // buffer
-                 eltype,  // eltype
-                 size,    // numel
-                 0,       // stride,
-                 0        // offset
-    );
+    NdArrayRef arr(buf,                                    // buffer
+                   eltype,                                 // eltype
+                   shape,                                  // shape
+                   std::vector<int64_t>(shape.size(), 0),  // strides
+                   0);
 
     DISPATCH_ALL_FIELDS(field, "pub2k.make_p", [&]() {
-      res.at<ring2k_t>(0) = static_cast<ring2k_t>(init);
+      arr.at<ring2k_t>(std::vector<int64_t>(shape.size(), 0)) =
+          static_cast<ring2k_t>(init);
     });
-    return res;
+    return Value(arr, DT_INVALID);
   }
 };
 
@@ -66,21 +66,16 @@ class Pub2kRandP : public Kernel {
   ce::CExpr comm() const override { return ce::Const(0); }
 
   void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<size_t>(0)));
+    ctx->setOutput(proc(ctx, ctx->getParam<Shape>(0)));
   }
 
-  static ArrayRef proc(KernelEvalContext* ctx, size_t size) {
-    SPU_TRACE_MPC_LEAF(ctx, size);
+  static Value proc(KernelEvalContext* ctx, const Shape& shape) {
     auto* prg_state = ctx->getState<PrgState>();
     const auto field = ctx->getState<Z2kState>()->getDefaultField();
 
-    // NOTE(junfeng): rand_p is heavily in unit tests and nbits has to been kept
-    // full.
-    //
-    // Reference: https://eprint.iacr.org/2019/599.pdf
-    // To make `msb based comparison` work, the safe range is
-    // [-2^(k-2), 2^(k-2))
-    return prg_state->genPubl(field, size).as(makeType<Pub2kTy>(field));
+    auto arr =
+        prg_state->genPubl(field, shape.numel()).as(makeType<Pub2kTy>(field));
+    return WrapValue(arr, shape);
   }
 };
 
@@ -93,7 +88,6 @@ class Pub2kNotP : public UnaryKernel {
   ce::CExpr comm() const override { return ce::Const(0); }
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in) const override {
-    SPU_TRACE_MPC_LEAF(ctx, in);
     const auto field = in.eltype().as<Ring2k>()->field();
     return ring_not(in).as(makeType<Pub2kTy>(field));
   }
@@ -109,7 +103,6 @@ class Pub2kEqualPP : public BinaryKernel {
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& x,
                 const ArrayRef& y) const override {
-    SPU_TRACE_MPC_LEAF(ctx, x, y);
     SPU_ENFORCE_EQ(x.eltype(), y.eltype());
     SPU_ENFORCE(x.eltype().isa<Pub2kTy>());
 
@@ -127,7 +120,6 @@ class Pub2kAddPP : public BinaryKernel {
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                 const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
     SPU_ENFORCE(lhs.eltype() == rhs.eltype());
     return ring_add(lhs, rhs).as(lhs.eltype());
   }
@@ -143,7 +135,6 @@ class Pub2kMulPP : public BinaryKernel {
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                 const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
     SPU_ENFORCE(lhs.eltype() == rhs.eltype());
     return ring_mul(lhs, rhs).as(lhs.eltype());
   }
@@ -160,7 +151,6 @@ class Pub2kMatMulPP : public MatmulKernel {
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                 const ArrayRef& rhs, size_t m, size_t n,
                 size_t k) const override {
-    SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
     SPU_ENFORCE(lhs.eltype() == rhs.eltype());
     return ring_mmul(lhs, rhs, m, n, k).as(lhs.eltype());
   }
@@ -176,7 +166,6 @@ class Pub2kAndPP : public BinaryKernel {
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                 const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
     SPU_ENFORCE(lhs.eltype() == rhs.eltype());
     return ring_and(lhs, rhs).as(lhs.eltype());
   }
@@ -192,7 +181,6 @@ class Pub2kXorPP : public BinaryKernel {
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                 const ArrayRef& rhs) const override {
-    SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
     SPU_ENFORCE(lhs.eltype() == rhs.eltype());
     return ring_xor(lhs, rhs).as(lhs.eltype());
   }
@@ -208,7 +196,6 @@ class Pub2kLShiftP : public ShiftKernel {
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
                 size_t bits) const override {
-    SPU_TRACE_MPC_LEAF(ctx, in, bits);
     return ring_lshift(in, bits).as(in.eltype());
   }
 };
@@ -223,7 +210,6 @@ class Pub2kRShiftP : public ShiftKernel {
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
                 size_t bits) const override {
-    SPU_TRACE_MPC_LEAF(ctx, in, bits);
     return ring_rshift(in, bits).as(in.eltype());
   }
 };
@@ -242,7 +228,6 @@ class Pub2kBitrevP : public BitrevKernel {
     SPU_ENFORCE(start <= end);
     SPU_ENFORCE(end <= SizeOf(field) * 8);
 
-    SPU_TRACE_MPC_LEAF(ctx, in, start, end);
     return ring_bitrev(in, start, end).as(in.eltype());
   }
 };
@@ -257,7 +242,6 @@ class Pub2kARShiftP : public ShiftKernel {
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
                 size_t bits) const override {
-    SPU_TRACE_MPC_LEAF(ctx, in, bits);
     return ring_arshift(in, bits).as(in.eltype());
   }
 };
@@ -272,7 +256,6 @@ class Pub2kTruncP : public ShiftKernel {
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
                 size_t bits) const override {
-    SPU_TRACE_MPC_LEAF(ctx, in, bits);
     // Rounding
     // AxB = (AxB >> 14) + ((AxB >> 13) & 1);
     // See
@@ -296,7 +279,6 @@ class Pub2kMsbP : public UnaryKernel {
   ce::CExpr comm() const override { return ce::Const(0); }
 
   ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in) const override {
-    SPU_TRACE_MPC_LEAF(ctx, in);
     return ring_rshift(in, in.elsize() * 8 - 1).as(in.eltype());
   }
 };
diff --git a/libspu/mpc/common/pub2k.h b/libspu/mpc/common/pub2k.h
index a387ed62..b94a1668 100644
--- a/libspu/mpc/common/pub2k.h
+++ b/libspu/mpc/common/pub2k.h
@@ -14,8 +14,8 @@
 
 #pragma once
 
+#include "libspu/core/object.h"
 #include "libspu/core/type.h"
-#include "libspu/mpc/object.h"
 
 namespace spu::mpc {
 
diff --git a/libspu/mpc/factory.cc b/libspu/mpc/factory.cc
index 42989e6d..d4bc8ea4 100644
--- a/libspu/mpc/factory.cc
+++ b/libspu/mpc/factory.cc
@@ -27,27 +27,26 @@
 
 namespace spu::mpc {
 
-std::unique_ptr<Object> Factory::CreateCompute(
-    const RuntimeConfig& conf,
-    const std::shared_ptr<yacl::link::Context>& lctx) {
-  switch (conf.protocol()) {
+void Factory::RegisterProtocol(
+    SPUContext* ctx, const std::shared_ptr<yacl::link::Context>& lctx) {
+  // TODO: support multi-protocols.
+  switch (ctx->config().protocol()) {
     case ProtocolKind::REF2K: {
-      return makeRef2kProtocol(conf, lctx);
+      return regRef2kProtocol(ctx, lctx);
     }
     case ProtocolKind::SEMI2K: {
-      return makeSemi2kProtocol(conf, lctx);
+      return regSemi2kProtocol(ctx, lctx);
     }
     case ProtocolKind::ABY3: {
-      return makeAby3Protocol(conf, lctx);
+      return regAby3Protocol(ctx, lctx);
     }
     case ProtocolKind::CHEETAH: {
-      return makeCheetahProtocol(conf, lctx);
+      return regCheetahProtocol(ctx, lctx);
     }
     default: {
-      SPU_THROW("Invalid protocol kind {}", conf.protocol());
+      SPU_THROW("Invalid protocol kind {}", ctx->config().protocol());
     }
   }
-  return nullptr;
 }
 
 std::unique_ptr<IoInterface> Factory::CreateIO(const RuntimeConfig& conf,
diff --git a/libspu/mpc/factory.h b/libspu/mpc/factory.h
index 0357c0c4..5583faa8 100644
--- a/libspu/mpc/factory.h
+++ b/libspu/mpc/factory.h
@@ -18,8 +18,8 @@
 
 #include "yacl/link/link.h"
 
+#include "libspu/core/context.h"
 #include "libspu/mpc/io_interface.h"
-#include "libspu/mpc/object.h"
 
 #include "libspu/spu.pb.h"
 
@@ -27,13 +27,12 @@ namespace spu::mpc {
 
 class Factory final {
  public:
-  // Create a computation context with given link.
+  // Add a protocol to a context.
   //
-  // @param kind, the protocol kind.
+  // @param config, a runtime config.
   // @param lctx, the inter party link context.
-  static std::unique_ptr<Object> CreateCompute(
-      const RuntimeConfig& conf,
-      const std::shared_ptr<yacl::link::Context>& lctx);
+  static void RegisterProtocol(
+      SPUContext* ctx, const std::shared_ptr<yacl::link::Context>& lctx);
 
   // Create a io context.
   //
diff --git a/libspu/mpc/kernel.h b/libspu/mpc/kernel.h
index b9e77ee1..207528b3 100644
--- a/libspu/mpc/kernel.h
+++ b/libspu/mpc/kernel.h
@@ -14,14 +14,32 @@
 
 #pragma once
 
-#include "libspu/mpc/object.h"
+#include "libspu/core/context.h"
+#include "libspu/core/prelude.h"
 
 namespace spu::mpc {
 
+class RandKernel : public Kernel {
+ public:
+  void evaluate(KernelEvalContext* ctx) const override {
+    const auto& shape = ctx->getParam<Shape>(0);
+
+    ArrayRef res = proc(ctx, shape.numel());
+
+    ctx->setOutput(WrapValue(res, shape));
+  }
+  virtual ArrayRef proc(KernelEvalContext* ctx, size_t size) const = 0;
+};
+
 class UnaryKernel : public Kernel {
  public:
   void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<ArrayRef>(0)));
+    const auto& in = ctx->getParam<Value>(0);
+    auto [arr, shape, dtype] = UnwrapValue(in);
+
+    ArrayRef res = proc(ctx, arr);
+
+    ctx->setOutput(WrapValue(res, shape));
   }
   virtual ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in) const = 0;
 };
@@ -29,8 +47,13 @@ class UnaryKernel : public Kernel {
 class ShiftKernel : public Kernel {
  public:
   void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(
-        proc(ctx, ctx->getParam<ArrayRef>(0), ctx->getParam<size_t>(1)));
+    const auto& in = ctx->getParam<Value>(0);
+    size_t bits = ctx->getParam<size_t>(1);
+    auto [arr, shape, dtype] = UnwrapValue(in);
+
+    ArrayRef res = proc(ctx, arr, bits);
+
+    ctx->setOutput(WrapValue(res, shape));
   }
   virtual ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
                         size_t bits) const = 0;
@@ -39,8 +62,18 @@ class ShiftKernel : public Kernel {
 class BinaryKernel : public Kernel {
  public:
   void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(
-        proc(ctx, ctx->getParam<ArrayRef>(0), ctx->getParam<ArrayRef>(1)));
+    const auto& lhs = ctx->getParam<Value>(0);
+    const auto& rhs = ctx->getParam<Value>(1);
+
+    SPU_ENFORCE(lhs.shape() == rhs.shape(), "shape mismatch {} {}", lhs.shape(),
+                rhs.shape());
+
+    auto [x, shape, dtype] = UnwrapValue(lhs);
+    auto [y, _, _1] = UnwrapValue(rhs);
+
+    auto z = proc(ctx, x, y);
+
+    ctx->setOutput(WrapValue(z, shape));
   }
   virtual ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                         const ArrayRef& rhs) const = 0;
@@ -49,10 +82,30 @@ class BinaryKernel : public Kernel {
 class MatmulKernel : public Kernel {
  public:
   void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<ArrayRef>(0),
-                        ctx->getParam<ArrayRef>(1), ctx->getParam<size_t>(2),
-                        ctx->getParam<size_t>(3), ctx->getParam<size_t>(4)));
+    const auto& lhs = ctx->getParam<Value>(0);
+    const auto& rhs = ctx->getParam<Value>(1);
+
+    // TODO: drop (m, n, k)
+    auto m = static_cast<int64_t>(ctx->getParam<size_t>(2));
+    auto n = static_cast<int64_t>(ctx->getParam<size_t>(3));
+    auto k = static_cast<int64_t>(ctx->getParam<size_t>(4));
+
+    // SPU_ENFORCE(lhs.shape().size() == 2 && rhs.shape().size() == 2 &&
+    //                 lhs.shape()[0] == m && lhs.shape()[1] == k &&
+    //                 rhs.shape()[0] == k && rhs.shape()[1] == n,
+    //             "invalid shape {} {}", lhs.shape(), rhs.shape());
+    SPU_ENFORCE(
+        calcNumel(lhs.shape()) == m * k && calcNumel(rhs.shape()) == k * n,
+        "invalid shape {} {}", lhs.shape(), rhs.shape());
+
+    auto [x, shape, dtype] = UnwrapValue(lhs);
+    auto [y, _, _1] = UnwrapValue(rhs);
+
+    auto z = proc(ctx, x, y, m, n, k);
+
+    ctx->setOutput(WrapValue(z, Shape{m, n}));
   }
+
   virtual ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& a,
                         const ArrayRef& b, size_t m, size_t n,
                         size_t k) const = 0;
@@ -61,12 +114,24 @@ class MatmulKernel : public Kernel {
 class Conv2DKernel : public Kernel {
  public:
   void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<ArrayRef>(0),
-                        ctx->getParam<ArrayRef>(1), ctx->getParam<size_t>(2),
-                        ctx->getParam<size_t>(3), ctx->getParam<size_t>(4),
-                        ctx->getParam<size_t>(5), ctx->getParam<size_t>(6),
-                        ctx->getParam<size_t>(7), ctx->getParam<size_t>(8),
-                        ctx->getParam<size_t>(9), ctx->getParam<size_t>(10)));
+    const auto& lhs = ctx->getParam<Value>(0);
+    const auto& rhs = ctx->getParam<Value>(1);
+    size_t N = ctx->getParam<size_t>(2);
+    size_t H = ctx->getParam<size_t>(3);
+    size_t W = ctx->getParam<size_t>(4);
+    size_t C = ctx->getParam<size_t>(5);
+    size_t O = ctx->getParam<size_t>(6);
+    size_t h = ctx->getParam<size_t>(7);
+    size_t w = ctx->getParam<size_t>(8);
+    size_t stride_h = ctx->getParam<size_t>(9);
+    size_t stride_w = ctx->getParam<size_t>(10);
+
+    auto [x, shape, dtype] = UnwrapValue(lhs);
+    auto [y, _, _1] = UnwrapValue(rhs);
+
+    auto z = proc(ctx, x, y, N, H, W, C, O, h, w, stride_h, stride_w);
+
+    ctx->setOutput(WrapValue(z, shape));
   }
   // tensor: NxHxWxC
   // filter: hxwxCxO
@@ -79,9 +144,17 @@ class Conv2DKernel : public Kernel {
 class BitrevKernel : public Kernel {
  public:
   void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<ArrayRef>(0),
-                        ctx->getParam<size_t>(1), ctx->getParam<size_t>(2)));
+    const auto& in = ctx->getParam<Value>(0);
+    size_t start = ctx->getParam<size_t>(1);
+    size_t end = ctx->getParam<size_t>(2);
+
+    auto [x, shape, dtype] = UnwrapValue(in);
+
+    auto z = proc(ctx, x, start, end);
+
+    ctx->setOutput(WrapValue(z, shape));
   }
+
   virtual ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
                         size_t start, size_t end) const = 0;
 };
@@ -112,8 +185,14 @@ class TruncAKernel : public ShiftKernel {
 class TruncAWithSignKernel : public Kernel {
  public:
   void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<ArrayRef>(0),
-                        ctx->getParam<size_t>(1), ctx->getParam<bool>(2)));
+    const auto& in = ctx->getParam<Value>(0);
+    size_t bits = ctx->getParam<size_t>(1);
+    bool positive = ctx->getParam<bool>(2);
+
+    auto [x, shape, dtype] = UnwrapValue(in);
+    auto z = proc(ctx, x, bits, positive);
+
+    ctx->setOutput(WrapValue(z, shape));
   }
 
   virtual bool hasMsbError() const = 0;
@@ -124,4 +203,35 @@ class TruncAWithSignKernel : public Kernel {
                         bool is_positive) const = 0;
 };
 
+class BitSplitKernel : public Kernel {
+ public:
+  void evaluate(KernelEvalContext* ctx) const override {
+    const auto& in = ctx->getParam<Value>(0);
+    size_t stride = ctx->getParam<size_t>(1);
+    auto [arr, shape, dtype] = UnwrapValue(in);
+
+    ArrayRef res = proc(ctx, arr, stride);
+
+    ctx->setOutput(WrapValue(res, shape));
+  }
+  virtual ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
+                        size_t stride) const = 0;
+};
+
+class CastTypeKernel : public Kernel {
+  void evaluate(KernelEvalContext* ctx) const override {
+    const auto& val = ctx->getParam<Value>(0);
+    const auto& to_type = ctx->getParam<Type>(1);
+
+    auto [arr, shape, dtype] = UnwrapValue(val);
+
+    ArrayRef res = proc(ctx, arr, to_type);
+
+    ctx->setOutput(WrapValue(res, shape));
+  }
+
+  virtual ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
+                        const Type& to_type) const = 0;
+};
+
 }  // namespace spu::mpc
diff --git a/libspu/mpc/object.h b/libspu/mpc/object.h
deleted file mode 100644
index 0bfde276..00000000
--- a/libspu/mpc/object.h
+++ /dev/null
@@ -1,266 +0,0 @@
-// Copyright 2021 Ant Group Co., Ltd.
-//
-// 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.
-
-#pragma once
-
-#include <map>
-#include <memory>
-#include <optional>
-#include <string>
-#include <variant>
-
-#include "libspu/core/array_ref.h"
-#include "libspu/core/prelude.h"
-#include "libspu/mpc/utils/cexpr.h"
-
-namespace spu::mpc {
-
-// TODO(jint) document the object model.
-// - KernelEvalContext: the calling convention
-// - Kernel: the dynamic member function.
-// - State: the dynamic member variable.
-// - Object: the dynamic binding object.
-
-// Helper class to instantiate kernel calls.
-template <typename CallerT>
-class EvaluationContext final {
-  // Please keep param types as less as possible.
-  using ParamType = std::variant<  //
-      std::optional<ArrayRef>,  // for kernel evaluation failure, give a change
-                                // for the caller to handle failed case.
-      ArrayRef,                 // value type
-      size_t,                   // represent size(mmul), shift_bits(shift)
-      bool,                     // binary flag
-      Type,                     // type of type
-      uint128_t                 // ring constant
-      >;
-
-  CallerT* caller_;
-  std::vector<ParamType> params_;
-  ParamType output_;
-
- public:
-  explicit EvaluationContext(CallerT* caller) : caller_(caller) {}
-
-  size_t numParams() const { return params_.size(); }
-
-  // Steal the output from this evaluation context.
-  //
-  // * usually called by kernel caller.
-  template <typename T = ArrayRef>
-  T&& stealOutput() {
-    return std::move(std::get<T>(output_));
-  }
-
-  // Bind an input to this evaluation context.
-  //
-  // * usually called by kernel caller.
-  template <typename T>
-  void bindParam(const T& in) {
-    params_.emplace_back(in);
-  }
-
-  // Get the caller's pointer.
-  //
-  // * usually called by kernel callee.
-  CallerT* caller() const {
-    if (auto caller = dynamic_cast<CallerT*>(caller_)) {
-      return caller;
-    }
-    SPU_THROW("cast failed");
-  }
-
-  // Get the i'th parameter.
-  //
-  // * usually called by kernel callee.
-  template <typename T>
-  const T& getParam(size_t pos) const {
-    SPU_ENFORCE(pos < params_.size(), "pos={} exceed num of inputs={}", pos,
-                params_.size());
-    return std::get<T>(params_[pos]);
-  }
-
-  // Set the output.
-  //
-  // * usually called by kernel callee.
-  template <typename T = ArrayRef>
-  void setOutput(T&& out) {
-    output_ = std::forward<T>(out);
-  }
-
-  // helper function, forward to caller.
-  template <typename StateT>
-  StateT* getState() {
-    return caller()->template getState<StateT>();
-  }
-
-  // helper function, forward to caller
-  bool hasKernel(std::string_view name) const {
-    return caller()->hasKernel(name);
-  }
-
-  // helper function, forward to caller
-  template <typename Ret = ArrayRef, typename... Args>
-  Ret call(std::string_view name, Args&&... args) {
-    return caller()->template call<Ret, Args...>(name,
-                                                 std::forward<Args>(args)...);
-  }
-};
-
-class Object;
-class Kernel {
- public:
-  enum class Kind {
-    // Indicate the kernel's complexity is static known.
-    //
-    // Typically, static kernel does not depend on runtime options, such like
-    // selecting different kernels according to different configs.
-    //
-    // By default, we should make kernels as 'Static' as possible.
-    Static,
-
-    // Indicate the kernel depends on runtime options, this kind of kernel is
-    // hard to analysis statically.
-    Dynamic,
-  };
-
-  virtual ~Kernel() = default;
-
-  virtual Kind kind() const { return Kind::Static; }
-
-  // Calculate number of comm rounds required for this kernel.
-  virtual ce::CExpr latency() const { return nullptr; }
-
-  // Calculate number of comm in bits.
-  virtual ce::CExpr comm() const { return nullptr; }
-
-  // Some kernel's communication can not be measured/implemented easily, this
-  // field tells the tolerance of theory value and implementation diff, in
-  // percentage. i.e.
-  //
-  // kernel's cost is infinitely close to 2.
-  //   comm(kernel) = 1 + 1/2 + 1/4 + 1/8 + ... = lim(2)
-  //
-  // in implementation, program may not handle bit, but instead use byte as the
-  // minimum unit, this make the error even larger.
-  virtual float getCommTolerance() const { return 0.0; }
-
-  // Evaluate this protocol within given context.
-  virtual void evaluate(EvaluationContext<Object>* ctx) const = 0;
-};
-
-class State {
- public:
-  virtual ~State() = default;
-
-  virtual bool hasLowCostFork() const { return false; }
-  virtual std::unique_ptr<State> fork();
-};
-
-// A (kernel) dynamic object dispatch a function to a kernel at runtime.
-//
-// Class that inherit from this class could do `dynamic binding`.
-class Object final {
-  std::map<std::string_view, std::shared_ptr<Kernel>> kernels_;
-  std::map<std::string_view, std::unique_ptr<State>> states_;
-
-  std::string id_;   // this object id.
-  std::string pid_;  // parent id.
-
-  int64_t child_counter_ = 0;
-
- public:
-  explicit Object(std::string id, std::string pid = "")
-      : id_(std::move(id)), pid_(std::move(pid)) {}
-
-  const std::string& id() const { return id_; }
-  const std::string& pid() const { return pid_; }
-
-  //
-  std::unique_ptr<Object> fork();
-
-  bool hasLowCostFork() const;
-
-  void regKernel(std::string_view name, std::unique_ptr<Kernel> kernel);
-
-  template <typename KernelT>
-  void regKernel() {
-    regKernel(KernelT::kBindName, std::make_unique<KernelT>());
-  }
-
-  template <typename KernelT>
-  void regKernel(std::string_view name) {
-    return regKernel(name, std::make_unique<KernelT>());
-  }
-
-  Kernel* getKernel(std::string_view name);
-  bool hasKernel(std::string_view name) const;
-
-  void addState(std::string_view name, std::unique_ptr<State> state) {
-    const auto& itr = states_.find(name);
-    SPU_ENFORCE(itr == states_.end(), "state={} already exist", name);
-    states_.emplace(name, std::move(state));
-  }
-
-  template <typename StateT, typename... Args>
-  void addState(Args&&... args) {
-    addState(StateT::kBindName,
-             std::make_unique<StateT>(std::forward<Args>(args)...));
-  }
-
-  template <typename StateT>
-  StateT* getState() {
-    const auto& itr = states_.find(StateT::kBindName);
-    SPU_ENFORCE(itr != states_.end(), "state={} not found", StateT::kBindName);
-    return dynamic_cast<StateT*>(itr->second.get());
-  }
-
-  //
-  std::vector<std::string_view> getKernelNames() const {
-    std::vector<std::string_view> names;
-    names.reserve(kernels_.size());
-    for (auto const& itr : kernels_) {
-      names.push_back(itr.first);
-    }
-    return names;
-  }
-
-  template <typename Ret>
-  Ret callImpl(Kernel* kernel, EvaluationContext<Object>* ctx) {
-    kernel->evaluate(ctx);
-    return ctx->stealOutput<Ret>();
-  }
-
-  template <typename Ret, typename First, typename... Args>
-  Ret callImpl(Kernel* kernel, EvaluationContext<Object>* ctx, First&& head,
-               Args&&... tail) {
-    ctx->bindParam(std::forward<First>(head));
-    if constexpr (sizeof...(Args) == 0) {
-      return callImpl<Ret>(kernel, ctx);
-    } else {
-      return callImpl<Ret>(kernel, ctx, std::forward<Args>(tail)...);
-    }
-  }
-
-  template <typename Ret = ArrayRef, typename... Args>
-  Ret call(std::string_view name, Args&&... args) {
-    Kernel* kernel = getKernel(name);
-    EvaluationContext<Object> ctx(this);
-    return callImpl<Ret>(kernel, &ctx, std::forward<Args>(args)...);
-  }
-};
-
-using KernelEvalContext = EvaluationContext<Object>;
-
-}  // namespace spu::mpc
diff --git a/libspu/mpc/ref2k/BUILD.bazel b/libspu/mpc/ref2k/BUILD.bazel
index 1f188580..96e2b56d 100644
--- a/libspu/mpc/ref2k/BUILD.bazel
+++ b/libspu/mpc/ref2k/BUILD.bazel
@@ -21,8 +21,8 @@ spu_cc_library(
     srcs = ["ref2k.cc"],
     hdrs = ["ref2k.h"],
     deps = [
+        "//libspu/core:context",
         "//libspu/mpc:io_interface",
-        "//libspu/mpc:object",
         "//libspu/mpc/common:prg_state",
         "//libspu/mpc/common:pub2k",
         "@yacl//yacl/link",
diff --git a/libspu/mpc/ref2k/ref2k.cc b/libspu/mpc/ref2k/ref2k.cc
index c709ba7c..23e75fd2 100644
--- a/libspu/mpc/ref2k/ref2k.cc
+++ b/libspu/mpc/ref2k/ref2k.cc
@@ -60,21 +60,19 @@ class Ref2kCommonTypeS : public Kernel {
   }
 };
 
-class Ref2kCastTypeS : public Kernel {
+class Ref2kCastTypeS : public CastTypeKernel {
  public:
   static constexpr char kBindName[] = "cast_type_s";
 
   Kind kind() const override { return Kind::Dynamic; }
 
-  void evaluate(KernelEvalContext* ctx) const override {
-    const auto& in = ctx->getParam<ArrayRef>(0);
-    const auto& to_type = ctx->getParam<Type>(1);
-
+  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
+                const Type& to_type) const override {
     SPU_TRACE_MPC_DISP(ctx, in, to_type);
     SPU_ENFORCE(in.eltype() == to_type,
                 "semi2k always use same bshare type, lhs={}, rhs={}",
                 in.eltype(), to_type);
-    ctx->setOutput(in);
+    return in;
   }
 };
 
@@ -104,7 +102,7 @@ class Ref2kS2P : public UnaryKernel {
   }
 };
 
-class Ref2kRandS : public Kernel {
+class Ref2kRandS : public RandKernel {
  public:
   static constexpr char kBindName[] = "rand_s";
 
@@ -112,12 +110,7 @@ class Ref2kRandS : public Kernel {
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<size_t>(0)));
-  }
-
-  static ArrayRef proc(KernelEvalContext* ctx, size_t size) {
-    SPU_TRACE_MPC_LEAF(ctx, size);
+  ArrayRef proc(KernelEvalContext* ctx, size_t size) const override {
     auto* state = ctx->getState<PrgState>();
     const auto field = ctx->getState<Z2kState>()->getDefaultField();
 
@@ -409,47 +402,52 @@ class Ref2kMsbS : public UnaryKernel {
 
 }  // namespace
 
-std::unique_ptr<Object> makeRef2kProtocol(
-    const RuntimeConfig& conf,
-    const std::shared_ptr<yacl::link::Context>& lctx) {
+void regRef2kProtocol(SPUContext* ctx,
+                      const std::shared_ptr<yacl::link::Context>& lctx) {
   registerTypes();
 
-  auto obj = std::make_unique<Object>("REF2K");
-
   // register random states & kernels.
-  obj->addState<PrgState>();
+  ctx->prot()->addState<PrgState>();
 
   // add Z2k state.
-  obj->addState<Z2kState>(conf.field());
+  ctx->prot()->addState<Z2kState>(ctx->config().field());
 
   // register public kernels.
-  regPub2kKernels(obj.get());
+  regPub2kKernels(ctx->prot());
 
   // register compute kernels
-  obj->regKernel<Ref2kCommonTypeS>();
-  obj->regKernel<Ref2kCastTypeS>();
-  obj->regKernel<Ref2kP2S>();
-  obj->regKernel<Ref2kS2P>();
-  obj->regKernel<Ref2kNotS>();
-  obj->regKernel<Ref2kAddSS>();
-  obj->regKernel<Ref2kAddSP>();
-  obj->regKernel<Ref2kMulSS>();
-  obj->regKernel<Ref2kMulSP>();
-  obj->regKernel<Ref2kMatMulSS>();
-  obj->regKernel<Ref2kMatMulSP>();
-  obj->regKernel<Ref2kAndSS>();
-  obj->regKernel<Ref2kAndSP>();
-  obj->regKernel<Ref2kXorSS>();
-  obj->regKernel<Ref2kXorSP>();
-  obj->regKernel<Ref2kLShiftS>();
-  obj->regKernel<Ref2kRShiftS>();
-  obj->regKernel<Ref2kBitrevS>();
-  obj->regKernel<Ref2kARShiftS>();
-  obj->regKernel<Ref2kTruncS>();
-  obj->regKernel<Ref2kMsbS>();
-  obj->regKernel<Ref2kRandS>();
-
-  return obj;
+  ctx->prot()->regKernel<Ref2kCommonTypeS>();
+  ctx->prot()->regKernel<Ref2kCastTypeS>();
+  ctx->prot()->regKernel<Ref2kP2S>();
+  ctx->prot()->regKernel<Ref2kS2P>();
+  ctx->prot()->regKernel<Ref2kNotS>();
+  ctx->prot()->regKernel<Ref2kAddSS>();
+  ctx->prot()->regKernel<Ref2kAddSP>();
+  ctx->prot()->regKernel<Ref2kMulSS>();
+  ctx->prot()->regKernel<Ref2kMulSP>();
+  ctx->prot()->regKernel<Ref2kMatMulSS>();
+  ctx->prot()->regKernel<Ref2kMatMulSP>();
+  ctx->prot()->regKernel<Ref2kAndSS>();
+  ctx->prot()->regKernel<Ref2kAndSP>();
+  ctx->prot()->regKernel<Ref2kXorSS>();
+  ctx->prot()->regKernel<Ref2kXorSP>();
+  ctx->prot()->regKernel<Ref2kLShiftS>();
+  ctx->prot()->regKernel<Ref2kRShiftS>();
+  ctx->prot()->regKernel<Ref2kBitrevS>();
+  ctx->prot()->regKernel<Ref2kARShiftS>();
+  ctx->prot()->regKernel<Ref2kTruncS>();
+  ctx->prot()->regKernel<Ref2kMsbS>();
+  ctx->prot()->regKernel<Ref2kRandS>();
+}
+
+std::unique_ptr<SPUContext> makeRef2kProtocol(
+    const RuntimeConfig& conf,
+    const std::shared_ptr<yacl::link::Context>& lctx) {
+  auto ctx = std::make_unique<SPUContext>(conf, lctx);
+
+  regRef2kProtocol(ctx.get(), lctx);
+
+  return ctx;
 }
 
 std::vector<ArrayRef> Ref2kIo::toShares(const ArrayRef& raw, Visibility vis,
diff --git a/libspu/mpc/ref2k/ref2k.h b/libspu/mpc/ref2k/ref2k.h
index d199fd41..c5d93ad1 100644
--- a/libspu/mpc/ref2k/ref2k.h
+++ b/libspu/mpc/ref2k/ref2k.h
@@ -16,8 +16,8 @@
 
 #include "yacl/link/link.h"
 
+#include "libspu/core/context.h"
 #include "libspu/mpc/io_interface.h"
-#include "libspu/mpc/object.h"
 
 namespace spu::mpc {
 
@@ -31,9 +31,11 @@ class Ref2kIo final : public BaseIo {
   ArrayRef fromShares(const std::vector<ArrayRef>& shares) const override;
 };
 
-std::unique_ptr<Object> makeRef2kProtocol(
+std::unique_ptr<SPUContext> makeRef2kProtocol(
     const RuntimeConfig& conf,
     const std::shared_ptr<yacl::link::Context>& lctx);
+void regRef2kProtocol(SPUContext* ctx,
+                      const std::shared_ptr<yacl::link::Context>& lctx);
 
 std::unique_ptr<Ref2kIo> makeRef2kIo(FieldType field, size_t npc);
 
diff --git a/libspu/mpc/semi2k/BUILD.bazel b/libspu/mpc/semi2k/BUILD.bazel
index 96b612f7..9e4bd190 100644
--- a/libspu/mpc/semi2k/BUILD.bazel
+++ b/libspu/mpc/semi2k/BUILD.bazel
@@ -40,8 +40,8 @@ spu_cc_library(
     deps = [
         ":state",
         ":type",
+        "//libspu/mpc:ab_api",
         "//libspu/mpc:kernel",
-        "//libspu/mpc/common:ab_api",
         "//libspu/mpc/common:communicator",
     ],
 )
@@ -54,9 +54,8 @@ spu_cc_library(
         ":state",
         ":type",
         "//libspu/core:vectorize",
+        "//libspu/mpc:ab_api",
         "//libspu/mpc:kernel",
-        "//libspu/mpc/common:ab_api",
-        "//libspu/mpc/common:ab_kernels",
         "//libspu/mpc/common:communicator",
     ],
 )
@@ -69,8 +68,8 @@ spu_cc_library(
         ":state",
         ":type",
         "//libspu/core:vectorize",
+        "//libspu/mpc:ab_api",
         "//libspu/mpc:kernel",
-        "//libspu/mpc/common:ab_api",
         "//libspu/mpc/common:communicator",
         "//libspu/mpc/utils:circuits",
         "//libspu/mpc/utils:ring_ops",
@@ -86,8 +85,6 @@ spu_cc_library(
         ":boolean",
         ":conversion",
         ":state",
-        "//libspu/mpc/common:ab_api",
-        "//libspu/mpc/common:ab_kernels",
         "//libspu/mpc/common:prg_state",
     ],
 )
@@ -97,8 +94,8 @@ spu_cc_test(
     srcs = ["protocol_test.cc"],
     deps = [
         ":protocol",
+        "//libspu/mpc:ab_api_test",
         "//libspu/mpc:api_test",
-        "//libspu/mpc/common:ab_api_test",
         "//libspu/mpc/semi2k/beaver/ttp_server:beaver_server",
     ],
 )
diff --git a/libspu/mpc/semi2k/arithmetic.cc b/libspu/mpc/semi2k/arithmetic.cc
index b01c1d49..8b0785c5 100644
--- a/libspu/mpc/semi2k/arithmetic.cc
+++ b/libspu/mpc/semi2k/arithmetic.cc
@@ -16,10 +16,9 @@
 
 #include <functional>
 
-#include "libspu/core/trace.h"
 #include "libspu/core/type_util.h"
 #include "libspu/core/vectorize.h"
-#include "libspu/mpc/common/ab_api.h"
+#include "libspu/mpc/ab_api.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/common/prg_state.h"
 #include "libspu/mpc/common/pub2k.h"
@@ -29,19 +28,7 @@
 
 namespace spu::mpc::semi2k {
 
-ArrayRef ZeroA::proc(KernelEvalContext* ctx, size_t size) {
-  SPU_TRACE_MPC_LEAF(ctx, size);
-
-  auto* prg_state = ctx->getState<PrgState>();
-  const auto field = ctx->getState<Z2kState>()->getDefaultField();
-
-  auto [r0, r1] = prg_state->genPrssPair(field, size);
-  return ring_sub(r0, r1).as(makeType<AShrTy>(field));
-}
-
-ArrayRef RandA::proc(KernelEvalContext* ctx, size_t size) {
-  SPU_TRACE_MPC_LEAF(ctx, size);
-
+ArrayRef RandA::proc(KernelEvalContext* ctx, size_t size) const {
   auto* prg_state = ctx->getState<PrgState>();
   const auto field = ctx->getState<Z2kState>()->getDefaultField();
 
@@ -57,12 +44,12 @@ ArrayRef RandA::proc(KernelEvalContext* ctx, size_t size) {
 }
 
 ArrayRef P2A::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   const auto field = in.eltype().as<Ring2k>()->field();
+  auto* prg_state = ctx->getState<PrgState>();
   auto* comm = ctx->getState<Communicator>();
 
-  auto x = zero_a(ctx->caller(), in.numel());
+  auto [r0, r1] = prg_state->genPrssPair(field, in.numel());
+  auto x = ring_sub(r0, r1).as(makeType<AShrTy>(field));
 
   if (comm->getRank() == 0) {
     ring_add_(x, in);
@@ -72,8 +59,6 @@ ArrayRef P2A::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 }
 
 ArrayRef A2P::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   const auto field = in.eltype().as<Ring2k>()->field();
   auto* comm = ctx->getState<Communicator>();
   auto out = comm->allReduce(ReduceOp::ADD, in, kBindName);
@@ -81,9 +66,7 @@ ArrayRef A2P::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 }
 
 ArrayRef NotA::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
-  auto* comm = ctx->caller()->getState<Communicator>();
+  auto* comm = ctx->getState<Communicator>();
 
   // First, let's show negate could be locally processed.
   //   let X = sum(Xi)     % M
@@ -114,8 +97,6 @@ ArrayRef NotA::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 ////////////////////////////////////////////////////////////////////
 ArrayRef AddAP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   SPU_ENFORCE(lhs.numel() == rhs.numel());
   auto* comm = ctx->getState<Communicator>();
 
@@ -128,8 +109,6 @@ ArrayRef AddAP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef AddAA::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   SPU_ENFORCE(lhs.numel() == rhs.numel());
   SPU_ENFORCE(lhs.eltype() == rhs.eltype());
 
@@ -141,18 +120,14 @@ ArrayRef AddAA::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 ////////////////////////////////////////////////////////////////////
 ArrayRef MulAP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   return ring_mul(lhs, rhs).as(lhs.eltype());
 }
 
 ArrayRef MulAA::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
-
   const auto field = lhs.eltype().as<Ring2k>()->field();
-  auto* comm = ctx->caller()->getState<Communicator>();
-  auto* beaver = ctx->caller()->getState<Semi2kState>()->beaver();
+  auto* comm = ctx->getState<Communicator>();
+  auto* beaver = ctx->getState<Semi2kState>()->beaver();
 
   auto res = ArrayRef(makeType<AShrTy>(field), lhs.numel());
   DISPATCH_ALL_FIELDS(field, kBindName, [&]() {
@@ -199,14 +174,11 @@ ArrayRef MulAA::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 ////////////////////////////////////////////////////////////////////
 ArrayRef MatMulAP::proc(KernelEvalContext* ctx, const ArrayRef& x,
                         const ArrayRef& y, size_t m, size_t n, size_t k) const {
-  SPU_TRACE_MPC_LEAF(ctx, x, y);
   return ring_mmul(x, y, m, n, k).as(x.eltype());
 }
 
 ArrayRef MatMulAA::proc(KernelEvalContext* ctx, const ArrayRef& x,
                         const ArrayRef& y, size_t m, size_t n, size_t k) const {
-  SPU_TRACE_MPC_LEAF(ctx, x, y);
-
   const auto field = x.eltype().as<Ring2k>()->field();
   auto* comm = ctx->getState<Communicator>();
   auto* beaver = ctx->getState<Semi2kState>()->beaver();
@@ -234,8 +206,6 @@ ArrayRef MatMulAA::proc(KernelEvalContext* ctx, const ArrayRef& x,
 
 ArrayRef LShiftA::proc(KernelEvalContext* ctx, const ArrayRef& in,
                        size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, bits);
-
   const auto field = in.eltype().as<Ring2k>()->field();
   bits %= SizeOf(field) * 8;
 
@@ -244,7 +214,6 @@ ArrayRef LShiftA::proc(KernelEvalContext* ctx, const ArrayRef& in,
 
 ArrayRef TruncA::proc(KernelEvalContext* ctx, const ArrayRef& x,
                       size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, x, bits);
   auto* comm = ctx->getState<Communicator>();
 
   // TODO: add truncation method to options.
@@ -274,7 +243,6 @@ ArrayRef TruncA::proc(KernelEvalContext* ctx, const ArrayRef& x,
 
 ArrayRef TruncAPr::proc(KernelEvalContext* ctx, const ArrayRef& in,
                         size_t bits) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, bits);
   auto* comm = ctx->getState<Communicator>();
   auto* beaver = ctx->getState<Semi2kState>()->beaver();
   const auto numel = in.numel();
diff --git a/libspu/mpc/semi2k/arithmetic.h b/libspu/mpc/semi2k/arithmetic.h
index df5439d9..88750ea2 100644
--- a/libspu/mpc/semi2k/arithmetic.h
+++ b/libspu/mpc/semi2k/arithmetic.h
@@ -15,26 +15,10 @@
 #pragma once
 
 #include "libspu/mpc/kernel.h"
-#include "libspu/mpc/utils/cexpr.h"
 
 namespace spu::mpc::semi2k {
 
-class ZeroA : public Kernel {
- public:
-  static constexpr char kBindName[] = "zero_a";
-
-  ce::CExpr latency() const override { return ce::Const(0); }
-
-  ce::CExpr comm() const override { return ce::Const(0); }
-
-  void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<size_t>(0)));
-  }
-
-  static ArrayRef proc(KernelEvalContext* ctx, size_t size);
-};
-
-class RandA : public Kernel {
+class RandA : public RandKernel {
  public:
   static constexpr char kBindName[] = "rand_a";
 
@@ -42,11 +26,7 @@ class RandA : public Kernel {
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<size_t>(0)));
-  }
-
-  static ArrayRef proc(KernelEvalContext* ctx, size_t size);
+  ArrayRef proc(KernelEvalContext* ctx, size_t size) const override;
 };
 
 class P2A : public UnaryKernel {
diff --git a/libspu/mpc/semi2k/boolean.cc b/libspu/mpc/semi2k/boolean.cc
index 5824aac6..33b40988 100644
--- a/libspu/mpc/semi2k/boolean.cc
+++ b/libspu/mpc/semi2k/boolean.cc
@@ -17,8 +17,7 @@
 #include <functional>
 
 #include "libspu/core/bit_utils.h"
-#include "libspu/core/trace.h"
-#include "libspu/mpc/common/ab_api.h"  // zero_b
+#include "libspu/mpc/ab_api.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/common/prg_state.h"
 #include "libspu/mpc/common/pub2k.h"
@@ -75,41 +74,21 @@ void CommonTypeB::evaluate(KernelEvalContext* ctx) const {
   const Type& lhs = ctx->getParam<Type>(0);
   const Type& rhs = ctx->getParam<Type>(1);
 
-  SPU_TRACE_MPC_DISP(ctx, lhs, rhs);
-
   SPU_ENFORCE(lhs == rhs, "semi2k always use same bshare type, lhs={}, rhs={}",
               lhs, rhs);
 
   ctx->setOutput(lhs);
 }
 
-void CastTypeB::evaluate(KernelEvalContext* ctx) const {
-  const auto& in = ctx->getParam<ArrayRef>(0);
-  const auto& to_type = ctx->getParam<Type>(1);
-
-  SPU_TRACE_MPC_DISP(ctx, in, to_type);
-
+ArrayRef CastTypeB::proc(KernelEvalContext* ctx, const ArrayRef& in,
+                         const Type& to_type) const {
   SPU_ENFORCE(in.eltype() == to_type,
               "semi2k always use same bshare type, lhs={}, rhs={}", in.eltype(),
               to_type);
-
-  ctx->setOutput(in);
-}
-
-ArrayRef ZeroB::proc(KernelEvalContext* ctx, size_t size) {
-  SPU_TRACE_MPC_LEAF(ctx, size);
-
-  // TODO: semantically, we should not use field for boolean share.
-  const auto field = ctx->getState<Z2kState>()->getDefaultField();
-  auto* prg_state = ctx->getState<PrgState>();
-  auto [r0, r1] = prg_state->genPrssPair(field, size);
-
-  return makeBShare(ring_xor(r0, r1), field, 0);
+  return in;
 }
 
 ArrayRef B2P::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   const auto field = in.eltype().as<Ring2k>()->field();
   auto* comm = ctx->getState<Communicator>();
   auto out = comm->allReduce(ReduceOp::XOR, in, kBindName);
@@ -117,11 +96,13 @@ ArrayRef B2P::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 }
 
 ArrayRef P2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
   const auto field = in.eltype().as<Ring2k>()->field();
+  auto* prg_state = ctx->getState<PrgState>();
+
   auto* comm = ctx->getState<Communicator>();
-  auto x = zero_b(ctx->caller(), in.numel());
+
+  auto [r0, r1] = prg_state->genPrssPair(field, in.numel());
+  auto x = ring_xor(r0, r1).as(makeType<BShrTy>(field, 0));
 
   if (comm->getRank() == 0) {
     ring_xor_(x, in);
@@ -132,7 +113,6 @@ ArrayRef P2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
 
 ArrayRef AndBP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
   SPU_ENFORCE(lhs.numel() == rhs.numel());
 
   const auto field = ctx->getState<Z2kState>()->getDefaultField();
@@ -153,7 +133,6 @@ ArrayRef AndBP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef AndBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
   SPU_ENFORCE(lhs.numel() == rhs.numel());
 
   auto* comm = ctx->getState<Communicator>();
@@ -213,7 +192,6 @@ ArrayRef AndBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef XorBP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
   SPU_ENFORCE(lhs.numel() == rhs.numel());
 
   auto* comm = ctx->getState<Communicator>();
@@ -230,7 +208,6 @@ ArrayRef XorBP::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef XorBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
                      const ArrayRef& rhs) const {
-  SPU_TRACE_MPC_LEAF(ctx, lhs, rhs);
   SPU_ENFORCE(lhs.numel() == rhs.numel());
 
   const auto field = ctx->getState<Z2kState>()->getDefaultField();
@@ -240,8 +217,6 @@ ArrayRef XorBB::proc(KernelEvalContext* ctx, const ArrayRef& lhs,
 
 ArrayRef LShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
                        size_t shift) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, shift);
-
   const auto field = in.eltype().as<Ring2k>()->field();
   shift %= SizeOf(field) * 8;
 
@@ -253,8 +228,6 @@ ArrayRef LShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
 
 ArrayRef RShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
                        size_t shift) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, shift);
-
   const auto field = in.eltype().as<Ring2k>()->field();
   shift %= SizeOf(field) * 8;
 
@@ -267,8 +240,6 @@ ArrayRef RShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
 
 ArrayRef ARShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
                         size_t shift) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, shift);
-
   const auto field = in.eltype().as<Ring2k>()->field();
   shift %= SizeOf(field) * 8;
 
@@ -279,7 +250,6 @@ ArrayRef ARShiftB::proc(KernelEvalContext* ctx, const ArrayRef& in,
 
 ArrayRef BitrevB::proc(KernelEvalContext* ctx, const ArrayRef& in, size_t start,
                        size_t end) const {
-  SPU_TRACE_MPC_LEAF(ctx, in, start, end);
   const auto field = in.eltype().as<Ring2k>()->field();
 
   SPU_ENFORCE(start <= end);
@@ -290,12 +260,8 @@ ArrayRef BitrevB::proc(KernelEvalContext* ctx, const ArrayRef& in, size_t start,
   return makeBShare(ring_bitrev(in, start, end), field, out_nbits);
 }
 
-void BitIntlB::evaluate(KernelEvalContext* ctx) const {
-  const auto& in = ctx->getParam<ArrayRef>(0);
-  const size_t stride = ctx->getParam<size_t>(1);
-
-  SPU_TRACE_MPC_LEAF(ctx, in, stride);
-
+ArrayRef BitIntlB::proc(KernelEvalContext* ctx, const ArrayRef& in,
+                        size_t stride) const {
   const auto field = in.eltype().as<Ring2k>()->field();
   const auto nbits = getNumBits(in);
   SPU_ENFORCE(absl::has_single_bit(nbits));
@@ -308,15 +274,11 @@ void BitIntlB::evaluate(KernelEvalContext* ctx) const {
     });
   });
 
-  ctx->setOutput(out);
+  return out;
 }
 
-void BitDeintlB::evaluate(KernelEvalContext* ctx) const {
-  const auto& in = ctx->getParam<ArrayRef>(0);
-  const size_t stride = ctx->getParam<size_t>(1);
-
-  SPU_TRACE_MPC_LEAF(ctx, in, stride);
-
+ArrayRef BitDeintlB::proc(KernelEvalContext* ctx, const ArrayRef& in,
+                          size_t stride) const {
   const auto field = in.eltype().as<Ring2k>()->field();
   const auto nbits = getNumBits(in);
   SPU_ENFORCE(absl::has_single_bit(nbits));
@@ -329,7 +291,7 @@ void BitDeintlB::evaluate(KernelEvalContext* ctx) const {
     });
   });
 
-  ctx->setOutput(out);
+  return out;
 }
 
 }  // namespace spu::mpc::semi2k
diff --git a/libspu/mpc/semi2k/boolean.h b/libspu/mpc/semi2k/boolean.h
index e584d313..6df523d5 100644
--- a/libspu/mpc/semi2k/boolean.h
+++ b/libspu/mpc/semi2k/boolean.h
@@ -15,7 +15,6 @@
 #pragma once
 
 #include "libspu/mpc/kernel.h"
-#include "libspu/mpc/utils/cexpr.h"
 
 namespace spu::mpc::semi2k {
 
@@ -28,28 +27,16 @@ class CommonTypeB : public Kernel {
   void evaluate(KernelEvalContext* ctx) const override;
 };
 
-class CastTypeB : public Kernel {
+class CastTypeB : public CastTypeKernel {
  public:
   static constexpr char kBindName[] = "cast_type_b";
 
-  Kind kind() const override { return Kind::Dynamic; }
-
-  void evaluate(KernelEvalContext* ctx) const override;
-};
-
-class ZeroB : public Kernel {
- public:
-  static constexpr char kBindName[] = "zero_b";
-
   ce::CExpr latency() const override { return ce::Const(0); }
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<size_t>(0)));
-  }
-
-  static ArrayRef proc(KernelEvalContext* ctx, size_t size);
+  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
+                const Type& to_type) const override;
 };
 
 class B2P : public UnaryKernel {
@@ -170,7 +157,7 @@ class BitrevB : public BitrevKernel {
                 size_t end) const override;
 };
 
-class BitIntlB : public Kernel {
+class BitIntlB : public BitSplitKernel {
  public:
   static constexpr char kBindName[] = "bitintl_b";
 
@@ -178,10 +165,11 @@ class BitIntlB : public Kernel {
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override;
+  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
+                size_t stride) const override;
 };
 
-class BitDeintlB : public Kernel {
+class BitDeintlB : public BitSplitKernel {
  public:
   static constexpr char kBindName[] = "bitdeintl_b";
 
@@ -189,7 +177,8 @@ class BitDeintlB : public Kernel {
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override;
+  ArrayRef proc(KernelEvalContext* ctx, const ArrayRef& in,
+                size_t stride) const override;
 };
 
 }  // namespace spu::mpc::semi2k
diff --git a/libspu/mpc/semi2k/conversion.cc b/libspu/mpc/semi2k/conversion.cc
index 95c1f533..956e2ba1 100644
--- a/libspu/mpc/semi2k/conversion.cc
+++ b/libspu/mpc/semi2k/conversion.cc
@@ -14,10 +14,8 @@
 
 #include "libspu/mpc/semi2k/conversion.h"
 
-#include "libspu/core/trace.h"
 #include "libspu/core/vectorize.h"
-#include "libspu/mpc/common/ab_api.h"
-#include "libspu/mpc/common/ab_kernels.h"
+#include "libspu/mpc/ab_api.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/common/prg_state.h"
 #include "libspu/mpc/semi2k/state.h"
@@ -26,16 +24,32 @@
 
 namespace spu::mpc::semi2k {
 
-ArrayRef A2B::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
-  SPU_TRACE_MPC_LEAF(ctx, x);
+static ArrayRef wrap_add_bb(SPUContext* ctx, const ArrayRef& x,
+                            const ArrayRef& y) {
+  SPU_ENFORCE(x.numel() == y.numel());
+  const Shape shape = {x.numel()};
+  auto [res, _s, _t] =
+      UnwrapValue(add_bb(ctx, WrapValue(x, shape), WrapValue(y, shape)));
+  return res;
+}
+
+static ArrayRef wrap_a2b(SPUContext* ctx, const ArrayRef& x) {
+  const Shape shape = {x.numel()};
+  auto [res, _s, _t] = UnwrapValue(a2b(ctx, WrapValue(x, shape)));
+  return res;
+}
 
+ArrayRef A2B::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
   const auto field = x.eltype().as<Ring2k>()->field();
   auto* comm = ctx->getState<Communicator>();
+  auto* prg_state = ctx->getState<PrgState>();
 
   std::vector<ArrayRef> bshrs;
   const auto bty = makeType<BShrTy>(field);
   for (size_t idx = 0; idx < comm->getWorldSize(); idx++) {
-    auto b = zero_b(ctx->caller(), x.numel());
+    auto [r0, r1] = prg_state->genPrssPair(field, x.numel());
+    auto b = ring_xor(r0, r1).as(bty);
+
     if (idx == comm->getRank()) {
       ring_xor_(b, x);
     }
@@ -44,14 +58,12 @@ ArrayRef A2B::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
 
   ArrayRef res = vectorizedReduce(bshrs.begin(), bshrs.end(),
                                   [&](const ArrayRef& xx, const ArrayRef& yy) {
-                                    return add_bb(ctx->caller(), xx, yy);
+                                    return wrap_add_bb(ctx->sctx(), xx, yy);
                                   });
-  return res.as(makeType<BShrTy>(field));
+  return res.as(bty);
 }
 
 ArrayRef B2A::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
-  SPU_TRACE_MPC_LEAF(ctx, x);
-
   const auto field = x.eltype().as<Ring2k>()->field();
   auto* comm = ctx->getState<Communicator>();
   auto* prg_state = ctx->getState<PrgState>();
@@ -60,11 +72,11 @@ ArrayRef B2A::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
   auto r_a = r_v.as(makeType<AShrTy>(field));
 
   // convert r to boolean share.
-  auto r_b = a2b(ctx->caller(), r_a);
+  auto r_b = wrap_a2b(ctx->sctx(), r_a);
 
   // evaluate adder circuit on x & r, and reveal x+r
-  auto x_plus_r =
-      comm->allReduce(ReduceOp::XOR, add_bb(ctx->caller(), x, r_b), kBindName);
+  auto x_plus_r = comm->allReduce(ReduceOp::XOR,
+                                  wrap_add_bb(ctx->sctx(), x, r_b), kBindName);
 
   // compute -r + (x+r)
   ring_neg_(r_a);
@@ -75,8 +87,6 @@ ArrayRef B2A::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
 }
 
 ArrayRef B2A_Randbit::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
-  SPU_TRACE_MPC_LEAF(ctx, x);
-
   const auto field = x.eltype().as<Ring2k>()->field();
   auto* comm = ctx->getState<Communicator>();
   auto* beaver = ctx->getState<Semi2kState>()->beaver();
@@ -130,10 +140,9 @@ ArrayRef B2A_Randbit::proc(KernelEvalContext* ctx, const ArrayRef& x) const {
 }
 
 ArrayRef MsbA2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
-  SPU_TRACE_MPC_LEAF(ctx, in);
-
-  const auto field = in.eltype().as<AShrTy>()->field();
+  const auto field = in.eltype().as<Ring2k>()->field();
   auto* comm = ctx->getState<Communicator>();
+  auto* prg_state = ctx->getState<PrgState>();
 
   // For if k > 2 parties does not collude with each other, then we can
   // construct two additive share and use carray out circuit directly.
@@ -143,7 +152,8 @@ ArrayRef MsbA2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
   std::vector<ArrayRef> bshrs;
   const auto bty = makeType<BShrTy>(field);
   for (size_t idx = 0; idx < comm->getWorldSize(); idx++) {
-    auto b = zero_b(ctx->caller(), in.numel());
+    auto [r0, r1] = prg_state->genPrssPair(field, in.numel());
+    auto b = ring_xor(r0, r1).as(bty);
     if (idx == comm->getRank()) {
       ring_xor_(b, in);
     }
@@ -156,9 +166,20 @@ ArrayRef MsbA2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
     k = 0;  // Empty matrix
   }
 
-  auto* obj = ctx->caller();
-  ArrayRef carry = common::carry_out(obj, bshrs[0], bshrs[1], k);
-  return xor_bb(obj, rshift_b(obj, xor_bb(obj, bshrs[0], bshrs[1]), k), carry);
+  auto* sctx = ctx->sctx();
+  const Shape shape = {in.numel()};
+  auto m = WrapValue(bshrs[0], shape);
+  auto n = WrapValue(bshrs[1], shape);
+  {
+    auto carry = carry_out(sctx, m, n, k);
+
+    // Compute the k'th bit.
+    //   (m^n)[k] ^ carry
+    auto msb = xor_bb(sctx, rshift_b(sctx, xor_bb(sctx, m, n), k), carry);
+
+    auto [msb_data, _shape, _dtype] = UnwrapValue(msb);
+    return msb_data;
+  }
 }
 
 }  // namespace spu::mpc::semi2k
diff --git a/libspu/mpc/semi2k/conversion.h b/libspu/mpc/semi2k/conversion.h
index 9831998c..01a12263 100644
--- a/libspu/mpc/semi2k/conversion.h
+++ b/libspu/mpc/semi2k/conversion.h
@@ -15,7 +15,6 @@
 #pragma once
 
 #include "libspu/mpc/kernel.h"
-#include "libspu/mpc/utils/cexpr.h"
 
 namespace spu::mpc::semi2k {
 
diff --git a/libspu/mpc/semi2k/protocol.cc b/libspu/mpc/semi2k/protocol.cc
index 3509a835..5055aff8 100644
--- a/libspu/mpc/semi2k/protocol.cc
+++ b/libspu/mpc/semi2k/protocol.cc
@@ -14,8 +14,6 @@
 
 #include "libspu/mpc/semi2k/protocol.h"
 
-#include "libspu/mpc/common/ab_api.h"
-#include "libspu/mpc/common/ab_kernels.h"
 #include "libspu/mpc/common/prg_state.h"
 #include "libspu/mpc/common/pub2k.h"
 #include "libspu/mpc/semi2k/arithmetic.h"
@@ -26,72 +24,71 @@
 
 namespace spu::mpc {
 
-std::unique_ptr<Object> makeSemi2kProtocol(
-    const RuntimeConfig& conf,
-    const std::shared_ptr<yacl::link::Context>& lctx) {
+void regSemi2kProtocol(SPUContext* ctx,
+                       const std::shared_ptr<yacl::link::Context>& lctx) {
   semi2k::registerTypes();
 
-  // FIXME: use same id for different rank
-  auto obj =
-      std::make_unique<Object>(fmt::format("{}-{}", lctx->Rank(), "SEMI2K"));
-
   // add communicator
-  obj->addState<Communicator>(lctx);
+  ctx->prot()->addState<Communicator>(lctx);
 
   // register random states & kernels.
-  obj->addState<PrgState>(lctx);
+  ctx->prot()->addState<PrgState>(lctx);
 
   // add Z2k state.
-  obj->addState<Z2kState>(conf.field());
+  ctx->prot()->addState<Z2kState>(ctx->config().field());
 
   // register public kernels.
-  regPub2kKernels(obj.get());
-
-  // register compute kernels
-  regABKernels(obj.get());
+  regPub2kKernels(ctx->prot());
 
   // register arithmetic & binary kernels
-  obj->addState<Semi2kState>(conf, lctx);
-  obj->regKernel<semi2k::ZeroA>();
-  obj->regKernel<semi2k::P2A>();
-  obj->regKernel<semi2k::A2P>();
-  obj->regKernel<semi2k::NotA>();
-  obj->regKernel<semi2k::AddAP>();
-  obj->regKernel<semi2k::AddAA>();
-  obj->regKernel<semi2k::MulAP>();
-  obj->regKernel<semi2k::MulAA>();
-  obj->regKernel<semi2k::MatMulAP>();
-  obj->regKernel<semi2k::MatMulAA>();
-  obj->regKernel<semi2k::LShiftA>();
-  // obj->regKernel<semi2k::TruncA>();
-  obj->regKernel<semi2k::TruncAPr>();
-
-  obj->regKernel<common::AddBB>();
-  obj->regKernel<semi2k::CommonTypeB>();
-  obj->regKernel<semi2k::CastTypeB>();
-  obj->regKernel<semi2k::ZeroB>();
-  obj->regKernel<semi2k::B2P>();
-  obj->regKernel<semi2k::P2B>();
-  obj->regKernel<semi2k::A2B>();
+  ctx->prot()->addState<Semi2kState>(ctx->config(), lctx);
+  ctx->prot()->regKernel<semi2k::P2A>();
+  ctx->prot()->regKernel<semi2k::A2P>();
+  ctx->prot()->regKernel<semi2k::NotA>();
+  ctx->prot()->regKernel<semi2k::AddAP>();
+  ctx->prot()->regKernel<semi2k::AddAA>();
+  ctx->prot()->regKernel<semi2k::MulAP>();
+  ctx->prot()->regKernel<semi2k::MulAA>();
+  ctx->prot()->regKernel<semi2k::MatMulAP>();
+  ctx->prot()->regKernel<semi2k::MatMulAA>();
+  ctx->prot()->regKernel<semi2k::LShiftA>();
+  // ctx->prot()->regKernel<semi2k::TruncA>();
+  ctx->prot()->regKernel<semi2k::TruncAPr>();
+
+  ctx->prot()->regKernel<semi2k::CommonTypeB>();
+  ctx->prot()->regKernel<semi2k::CastTypeB>();
+  ctx->prot()->regKernel<semi2k::B2P>();
+  ctx->prot()->regKernel<semi2k::P2B>();
+  ctx->prot()->regKernel<semi2k::A2B>();
 
   if (lctx->WorldSize() == 2) {
-    obj->regKernel<semi2k::MsbA2B>();
+    ctx->prot()->regKernel<semi2k::MsbA2B>();
   }
-  // obj->regKernel<semi2k::B2A>();
-  obj->regKernel<semi2k::B2A_Randbit>();
-  obj->regKernel<semi2k::AndBP>();
-  obj->regKernel<semi2k::AndBB>();
-  obj->regKernel<semi2k::XorBP>();
-  obj->regKernel<semi2k::XorBB>();
-  obj->regKernel<semi2k::LShiftB>();
-  obj->regKernel<semi2k::RShiftB>();
-  obj->regKernel<semi2k::ARShiftB>();
-  obj->regKernel<semi2k::BitrevB>();
-  obj->regKernel<semi2k::BitIntlB>();
-  obj->regKernel<semi2k::BitDeintlB>();
-  obj->regKernel<semi2k::RandA>();
-
-  return obj;
+  // ctx->prot()->regKernel<semi2k::B2A>();
+  ctx->prot()->regKernel<semi2k::B2A_Randbit>();
+  ctx->prot()->regKernel<semi2k::AndBP>();
+  ctx->prot()->regKernel<semi2k::AndBB>();
+  ctx->prot()->regKernel<semi2k::XorBP>();
+  ctx->prot()->regKernel<semi2k::XorBB>();
+  ctx->prot()->regKernel<semi2k::LShiftB>();
+  ctx->prot()->regKernel<semi2k::RShiftB>();
+  ctx->prot()->regKernel<semi2k::ARShiftB>();
+  ctx->prot()->regKernel<semi2k::BitrevB>();
+  ctx->prot()->regKernel<semi2k::BitIntlB>();
+  ctx->prot()->regKernel<semi2k::BitDeintlB>();
+  ctx->prot()->regKernel<semi2k::RandA>();
+}
+
+std::unique_ptr<SPUContext> makeSemi2kProtocol(
+    const RuntimeConfig& conf,
+    const std::shared_ptr<yacl::link::Context>& lctx) {
+  semi2k::registerTypes();
+
+  auto ctx = std::make_unique<SPUContext>(conf, lctx);
+
+  regSemi2kProtocol(ctx.get(), lctx);
+
+  return ctx;
 }
 
 }  // namespace spu::mpc
diff --git a/libspu/mpc/semi2k/protocol.h b/libspu/mpc/semi2k/protocol.h
index b1161c87..c27917d4 100644
--- a/libspu/mpc/semi2k/protocol.h
+++ b/libspu/mpc/semi2k/protocol.h
@@ -16,11 +16,14 @@
 
 #include "yacl/link/link.h"
 
-#include "libspu/mpc/object.h"
+#include "libspu/core/context.h"
 
 namespace spu::mpc {
 
-std::unique_ptr<Object> makeSemi2kProtocol(
+void regSemi2kProtocol(SPUContext* ctx,
+                       const std::shared_ptr<yacl::link::Context>& lctx);
+
+std::unique_ptr<SPUContext> makeSemi2kProtocol(
     const RuntimeConfig& conf,
     const std::shared_ptr<yacl::link::Context>& lctx);
 
diff --git a/libspu/mpc/spdz2k/BUILD.bazel b/libspu/mpc/spdz2k/BUILD.bazel
index 3d63960f..3cb550a9 100644
--- a/libspu/mpc/spdz2k/BUILD.bazel
+++ b/libspu/mpc/spdz2k/BUILD.bazel
@@ -32,7 +32,7 @@ spu_cc_library(
         ":arithmetic",
         ":state",
         ":value",
-        "//libspu/mpc/common:ab_api",
+        "//libspu/core:context",
         "//libspu/mpc/common:prg_state",
     ],
 )
@@ -64,8 +64,8 @@ spu_cc_library(
         ":type",
         ":value",
         "//libspu/core:vectorize",
+        "//libspu/mpc:ab_api",
         "//libspu/mpc:kernel",
-        "//libspu/mpc/common:ab_api",
         "//libspu/mpc/common:communicator",
         "//libspu/mpc/utils:circuits",
         "//libspu/mpc/utils:ring_ops",
@@ -141,8 +141,8 @@ spu_cc_library(
     srcs = ["abprotocol_spdz2k_test.cc"],
     hdrs = ["abprotocol_spdz2k_test.h"],
     deps = [
-        "//libspu/mpc:object",
-        "//libspu/mpc/common:ab_api",
+        "//libspu/mpc:ab_api",
+        "//libspu/mpc:api",
         "//libspu/mpc/common:communicator",
         "//libspu/mpc/common:pub2k",
         "//libspu/mpc/utils:simulate",
diff --git a/libspu/mpc/spdz2k/arithmetic.cc b/libspu/mpc/spdz2k/arithmetic.cc
index ea713891..376499b1 100644
--- a/libspu/mpc/spdz2k/arithmetic.cc
+++ b/libspu/mpc/spdz2k/arithmetic.cc
@@ -26,7 +26,7 @@
 #include "libspu/core/array_ref.h"
 #include "libspu/core/trace.h"
 #include "libspu/core/vectorize.h"
-#include "libspu/mpc/common/ab_api.h"
+#include "libspu/mpc/ab_api.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/common/prg_state.h"
 #include "libspu/mpc/common/pub2k.h"
@@ -37,10 +37,9 @@
 #include "libspu/mpc/utils/ring_ops.h"
 
 namespace spu::mpc::spdz2k {
+namespace {
 
-ArrayRef ZeroA::proc(KernelEvalContext* ctx, size_t size) {
-  SPU_TRACE_MPC_LEAF(ctx, size);
-
+ArrayRef zero_a_impl(KernelEvalContext* ctx, size_t size) {
   auto* prg_state = ctx->getState<PrgState>();
   const auto field = ctx->getState<Z2kState>()->getDefaultField();
 
@@ -59,7 +58,9 @@ ArrayRef ZeroA::proc(KernelEvalContext* ctx, size_t size) {
   return makeAShare(x, x_mac, field);
 }
 
-ArrayRef RandA::proc(KernelEvalContext* ctx, size_t size) {
+}  // namespace
+
+ArrayRef RandA::proc(KernelEvalContext* ctx, size_t size) const {
   SPU_TRACE_MPC_LEAF(ctx, size);
   SPU_THROW("NotImplemented");
 }
@@ -70,7 +71,7 @@ ArrayRef P2A::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
   auto* comm = ctx->getState<Communicator>();
   const auto key = ctx->getState<Spdz2kState>()->key();
 
-  auto res = zero_a(ctx->caller(), in.numel());
+  auto res = zero_a_impl(ctx, in.numel());
   auto z = getValueShare(res);
   auto z_mac = getMacShare(res);
 
@@ -231,6 +232,21 @@ bool SingleCheck(KernelEvalContext* ctx, const ArrayRef& in) {
   return ret;
 }
 
+static ArrayRef wrap_lshift_a(SPUContext* ctx, const ArrayRef& x, size_t k) {
+  const Shape shape = {x.numel()};
+  auto [res, _s, _t] = UnwrapValue(lshift_a(ctx, WrapValue(x, shape), k));
+  return res;
+}
+
+static ArrayRef wrap_add_aa(SPUContext* ctx, const ArrayRef& x,
+                            const ArrayRef& y) {
+  SPU_ENFORCE(x.numel() == y.numel());
+  const Shape shape = {x.numel()};
+  auto [res, _s, _t] =
+      UnwrapValue(add_aa(ctx, WrapValue(x, shape), WrapValue(y, shape)));
+  return res;
+}
+
 // Refer to:
 // Procedure BatchCheck, 3.2 Batch MAC Checking with Random Linear
 // Combinations, SPDZ2k: Efficient MPC mod 2k for Dishonest Majority
@@ -262,8 +278,8 @@ bool BatchCheck(KernelEvalContext* ctx, const std::vector<ArrayRef>& ins) {
     auto rmac = makeAShare(r, r_mac, field);
 
     // 2. [x_hat] = [x] + 2^k * [r]
-    auto l_rmac = lshift_a(ctx->caller(), rmac, k);
-    auto _in = add_aa(ctx->caller(), in, l_rmac);
+    auto l_rmac = wrap_lshift_a(ctx->sctx(), rmac, k);
+    auto _in = wrap_add_aa(ctx->sctx(), in, l_rmac);
 
     const auto& x_hat = getValueShare(_in);
     x_hat_v.emplace_back(x_hat);
diff --git a/libspu/mpc/spdz2k/arithmetic.h b/libspu/mpc/spdz2k/arithmetic.h
index 8152f0bd..8fa69ca0 100644
--- a/libspu/mpc/spdz2k/arithmetic.h
+++ b/libspu/mpc/spdz2k/arithmetic.h
@@ -15,26 +15,10 @@
 #pragma once
 
 #include "libspu/mpc/kernel.h"
-#include "libspu/mpc/utils/cexpr.h"
 
 namespace spu::mpc::spdz2k {
 
-class ZeroA : public Kernel {
- public:
-  static constexpr char kBindName[] = "zero_a";
-
-  ce::CExpr latency() const override { return ce::Const(0); }
-
-  ce::CExpr comm() const override { return ce::Const(0); }
-
-  void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<size_t>(0)));
-  }
-
-  static ArrayRef proc(KernelEvalContext* ctx, size_t size);
-};
-
-class RandA : public Kernel {
+class RandA : public RandKernel {
  public:
   static constexpr char kBindName[] = "rand_a";
 
@@ -42,11 +26,7 @@ class RandA : public Kernel {
 
   ce::CExpr comm() const override { return ce::Const(0); }
 
-  void evaluate(KernelEvalContext* ctx) const override {
-    ctx->setOutput(proc(ctx, ctx->getParam<size_t>(0)));
-  }
-
-  static ArrayRef proc(KernelEvalContext* ctx, size_t size);
+  ArrayRef proc(KernelEvalContext* ctx, size_t size) const override;
 };
 
 class P2A : public UnaryKernel {
diff --git a/libspu/mpc/spdz2k/protocol.cc b/libspu/mpc/spdz2k/protocol.cc
index d8599434..f9714717 100644
--- a/libspu/mpc/spdz2k/protocol.cc
+++ b/libspu/mpc/spdz2k/protocol.cc
@@ -14,58 +14,56 @@
 
 #include "libspu/mpc/spdz2k/protocol.h"
 
-#include "libspu/mpc/common/ab_api.h"
+#include "libspu/core/context.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/common/prg_state.h"
 #include "libspu/mpc/common/pub2k.h"
-#include "libspu/mpc/object.h"
 #include "libspu/mpc/spdz2k/arithmetic.h"
 #include "libspu/mpc/spdz2k/state.h"
 #include "libspu/mpc/spdz2k/type.h"
 
 namespace spu::mpc {
 
-std::unique_ptr<Object> makeSpdz2kProtocol(
-    const RuntimeConfig& conf,
-    const std::shared_ptr<yacl::link::Context>& lctx) {
+void regSpdz2kProtocol(SPUContext* ctx,
+                       const std::shared_ptr<yacl::link::Context>& lctx) {
   spdz2k::registerTypes();
 
-  // FIXME: use same id for different rank
-  auto obj =
-      std::make_unique<Object>(fmt::format("{}-{}", lctx->Rank(), "SPDZ2K"));
-
   // add communicator
-  obj->addState<Communicator>(lctx);
+  ctx->prot()->addState<Communicator>(lctx);
 
   // register random states & kernels.
-  obj->addState<PrgState>(lctx);
+  ctx->prot()->addState<PrgState>(lctx);
 
   // add Z2k state.
-  obj->addState<Z2kState>(conf.field());
+  ctx->prot()->addState<Z2kState>(ctx->config().field());
 
   // register public kernels.
-  regPub2kKernels(obj.get());
-
-  // register compute kernels
-  regABKernels(obj.get());
+  regPub2kKernels(ctx->prot());
 
   // register arithmetic kernels
-  obj->addState<Spdz2kState>(lctx);
-  obj->regKernel<spdz2k::ZeroA>();
-  obj->regKernel<spdz2k::P2A>();
-  obj->regKernel<spdz2k::A2P>();
-  obj->regKernel<spdz2k::NotA>();
-  obj->regKernel<spdz2k::AddAP>();
-  obj->regKernel<spdz2k::AddAA>();
-  obj->regKernel<spdz2k::MulAP>();
-  obj->regKernel<spdz2k::MulAA>();
-  obj->regKernel<spdz2k::MatMulAP>();
-  obj->regKernel<spdz2k::MatMulAA>();
-  obj->regKernel<spdz2k::LShiftA>();
-  obj->regKernel<spdz2k::TruncA>();
-  obj->regKernel<spdz2k::RandA>();
+  ctx->prot()->addState<Spdz2kState>(lctx);
+  ctx->prot()->regKernel<spdz2k::P2A>();
+  ctx->prot()->regKernel<spdz2k::A2P>();
+  ctx->prot()->regKernel<spdz2k::NotA>();
+  ctx->prot()->regKernel<spdz2k::AddAP>();
+  ctx->prot()->regKernel<spdz2k::AddAA>();
+  ctx->prot()->regKernel<spdz2k::MulAP>();
+  ctx->prot()->regKernel<spdz2k::MulAA>();
+  ctx->prot()->regKernel<spdz2k::MatMulAP>();
+  ctx->prot()->regKernel<spdz2k::MatMulAA>();
+  ctx->prot()->regKernel<spdz2k::LShiftA>();
+  ctx->prot()->regKernel<spdz2k::TruncA>();
+  ctx->prot()->regKernel<spdz2k::RandA>();
+}
+
+std::unique_ptr<SPUContext> makeSpdz2kProtocol(
+    const RuntimeConfig& conf,
+    const std::shared_ptr<yacl::link::Context>& lctx) {
+  auto ctx = std::make_unique<SPUContext>(conf, lctx);
+
+  regSpdz2kProtocol(ctx.get(), lctx);
 
-  return obj;
+  return ctx;
 }
 
 }  // namespace spu::mpc
diff --git a/libspu/mpc/spdz2k/protocol.h b/libspu/mpc/spdz2k/protocol.h
index 17495c49..31dcf0e4 100644
--- a/libspu/mpc/spdz2k/protocol.h
+++ b/libspu/mpc/spdz2k/protocol.h
@@ -16,11 +16,14 @@
 
 #include "yacl/link/link.h"
 
-#include "libspu/mpc/object.h"
+#include "libspu/core/context.h"
 
 namespace spu::mpc {
 
-std::unique_ptr<Object> makeSpdz2kProtocol(
+void regSpdz2kProtocol(SPUContext* ctx,
+                       const std::shared_ptr<yacl::link::Context>& lctx);
+
+std::unique_ptr<SPUContext> makeSpdz2kProtocol(
     const RuntimeConfig& conf,
     const std::shared_ptr<yacl::link::Context>& lctx);
 
diff --git a/libspu/mpc/spdz2k/state.h b/libspu/mpc/spdz2k/state.h
index e663212b..64438b96 100644
--- a/libspu/mpc/spdz2k/state.h
+++ b/libspu/mpc/spdz2k/state.h
@@ -21,8 +21,8 @@
 #include "yacl/link/link.h"
 
 #include "libspu/core/array_ref.h"
+#include "libspu/core/object.h"
 #include "libspu/mpc/common/communicator.h"
-#include "libspu/mpc/object.h"
 #include "libspu/mpc/spdz2k/beaver/beaver_tfp.h"
 #include "libspu/mpc/spdz2k/commitment.h"
 
diff --git a/libspu/mpc/tools/BUILD.bazel b/libspu/mpc/tools/BUILD.bazel
index 087feac7..d025c0bb 100644
--- a/libspu/mpc/tools/BUILD.bazel
+++ b/libspu/mpc/tools/BUILD.bazel
@@ -35,9 +35,10 @@ spu_cc_binary(
     srcs = ["complexity.cc"],
     deps = [
         ":complexity_cc_proto",
+        "//libspu/core:context",
+        "//libspu/mpc:api",
         "//libspu/mpc:factory",
         "//libspu/mpc:kernel",
-        "//libspu/mpc:object",
         "//libspu/mpc/utils:simulate",
         "@llvm-project//llvm:Support",
     ],
@@ -50,8 +51,8 @@ spu_cc_binary(
         "benchmark.h",
     ],
     deps = [
+        "//libspu/core:context",
         "//libspu/mpc:api",
-        "//libspu/mpc:object",
         "//libspu/mpc/aby3",
         "//libspu/mpc/cheetah",
         "//libspu/mpc/common:communicator",
diff --git a/libspu/mpc/tools/benchmark.cc b/libspu/mpc/tools/benchmark.cc
index 901060ce..66a1e2dc 100644
--- a/libspu/mpc/tools/benchmark.cc
+++ b/libspu/mpc/tools/benchmark.cc
@@ -173,9 +173,7 @@ DEFINE_BENCHMARK(BenchMMulSP, MatrixSizeArgs);
 DEFINE_BENCHMARK(BenchMMulSS, MatrixSizeArgs);
 
 DEFINE_BENCHMARK(BenchRandA, NumelArgs);
-DEFINE_BENCHMARK(BenchZeroA, NumelArgs);
 DEFINE_BENCHMARK(BenchRandB, NumelArgs);
-DEFINE_BENCHMARK(BenchZeroB, NumelArgs);
 DEFINE_BENCHMARK(BenchP2A, NumelArgs);
 DEFINE_BENCHMARK(BenchA2P, NumelArgs);
 DEFINE_BENCHMARK(BenchMsbA2b, NumelArgs);
diff --git a/libspu/mpc/tools/benchmark.h b/libspu/mpc/tools/benchmark.h
index ee57fd89..7201db3b 100644
--- a/libspu/mpc/tools/benchmark.h
+++ b/libspu/mpc/tools/benchmark.h
@@ -20,11 +20,11 @@
 #include "yacl/link/link.h"
 
 #include "libspu/core/bit_utils.h"
+#include "libspu/core/context.h"
 #include "libspu/core/shape_util.h"  // calcNumel
+#include "libspu/mpc/ab_api.h"
 #include "libspu/mpc/api.h"
-#include "libspu/mpc/common/ab_api.h"
 #include "libspu/mpc/common/communicator.h"
-#include "libspu/mpc/object.h"
 #include "libspu/mpc/utils/ring_ops.h"
 #include "libspu/mpc/utils/simulate.h"
 
@@ -37,7 +37,7 @@ struct ParamEntry {
       : name(std::move(n)), value(std::move(v)) {}
 };
 
-using CreateComputeFn = std::function<std::unique_ptr<Object>(
+using CreateComputeFn = std::function<std::unique_ptr<SPUContext>(
     const RuntimeConfig& conf,
     const std::shared_ptr<yacl::link::Context>& lctx)>;
 
@@ -48,7 +48,7 @@ class BenchConfig {
   inline static uint32_t bench_npc = 0;
   inline static std::string bench_mode = "standalone";
   inline static std::string bench_parties = {};
-  inline static std::vector<int64_t> bench_numel_range = {1u << 10, 1u << 20};
+  inline static std::vector<int64_t> bench_numel_range = {1U << 10, 1U << 20};
   inline static std::vector<int64_t> bench_shift_range = {2};
   inline static std::vector<int64_t> bench_matrix_range = {10, 100};
   inline static std::vector<int64_t> bench_field_range = {FieldType::FM64,
@@ -120,59 +120,59 @@ void MPCBenchMark(benchmark::State& state) {
   }
 }
 
-#define MPC_BENCH_DEFINE(CLASS, DATA, OP, ...)        \
-  class CLASS : public DATA {                         \
-   public:                                            \
-    using DATA::DATA;                                 \
-    static inline std::string op_name = #OP;          \
-    ArrayRef Exec() { return OP(obj_, __VA_ARGS__); } \
+#define MPC_BENCH_DEFINE(CLASS, DATA, OP, ...)     \
+  class CLASS : public DATA {                      \
+   public:                                         \
+    using DATA::DATA;                              \
+    static inline std::string op_name = #OP;       \
+    Value Exec() { return OP(obj_, __VA_ARGS__); } \
   };
 
 template <size_t P = 0, size_t S = 0, size_t A = 0, size_t B = 0, size_t MP = 0,
           size_t MS = 0, size_t MA = 0, size_t MB = 0, size_t B1 = 0>
 class OpData {
  protected:
-  Object* obj_{nullptr};
+  SPUContext* obj_{nullptr};
   benchmark::State& state;
-  std::array<ArrayRef, P> ps;
-  std::array<ArrayRef, S> ss;
-  std::array<ArrayRef, A> as;
-  std::array<ArrayRef, B> bs;
-  std::array<ArrayRef, MP> mps;
-  std::array<ArrayRef, MS> mss;
-  std::array<ArrayRef, MA> mas;
-  std::array<ArrayRef, MB> mbs;
-  std::array<ArrayRef, B1> b1s;
+  std::array<Value, P> ps;
+  std::array<Value, S> ss;
+  std::array<Value, A> as;
+  std::array<Value, B> bs;
+  std::array<Value, MP> mps;
+  std::array<Value, MS> mss;
+  std::array<Value, MA> mas;
+  std::array<Value, MB> mbs;
+  std::array<Value, B1> b1s;
 
  public:
-  OpData(Object* obj, benchmark::State& st) : obj_(obj), state(st) {
+  OpData(SPUContext* obj, benchmark::State& st) : obj_(obj), state(st) {
     for (auto& p : ps) {
-      p = rand_p(obj_, state.range(1));
+      p = rand_p(obj_, Shape{state.range(1)});
     }
     for (auto& s : ss) {
-      s = p2s(obj_, rand_p(obj, state.range(1)));
+      s = p2s(obj_, rand_p(obj, Shape{state.range(1)}));
     }
     for (auto& a : as) {
-      a = p2a(obj_, rand_p(obj_, state.range(1)));
+      a = p2a(obj_, rand_p(obj_, Shape{state.range(1)}));
     }
     for (auto& b : bs) {
-      b = p2b(obj_, rand_p(obj_, state.range(1)));
+      b = p2b(obj_, rand_p(obj_, Shape{state.range(1)}));
     }
     auto matrix_size = calcNumel({state.range(1), state.range(1)});
     for (auto& mp : mps) {
-      mp = rand_p(obj_, matrix_size);
+      mp = rand_p(obj_, Shape{matrix_size});
     }
     for (auto& ms : mss) {
-      ms = p2s(obj_, rand_p(obj_, matrix_size));
+      ms = p2s(obj_, rand_p(obj_, Shape{matrix_size}));
     }
     for (auto& ma : mas) {
-      ma = p2a(obj_, rand_p(obj_, matrix_size));
+      ma = p2a(obj_, rand_p(obj_, Shape{matrix_size}));
     }
     for (auto& mb : mbs) {
-      mb = p2b(obj_, rand_p(obj_, matrix_size));
+      mb = p2b(obj_, rand_p(obj_, Shape{matrix_size}));
     }
     for (auto& b1 : b1s) {
-      b1 = p2b(obj_, rand_p(obj_, state.range(1)));
+      b1 = p2b(obj_, rand_p(obj_, Shape{state.range(1)}));
       b1 = lshift_b(obj_, b1,
                     SizeOf(static_cast<FieldType>(state.range(0))) * 8 - 1);
       b1 = rshift_b(obj_, b1,
@@ -223,10 +223,8 @@ MPC_BENCH_DEFINE(BenchMMulSP, OpData1MS1MP, mmul_sp, mss[0], mps[0],
 MPC_BENCH_DEFINE(BenchMMulSS, OpData2MS, mmul_ss, mss[0], mss[1],
                  state.range(1), state.range(1), state.range(1))
 
-MPC_BENCH_DEFINE(BenchRandA, OpDataBasic, rand_a, state.range(1))
-MPC_BENCH_DEFINE(BenchZeroA, OpDataBasic, zero_a, state.range(1))
-MPC_BENCH_DEFINE(BenchRandB, OpDataBasic, rand_b, state.range(1))
-MPC_BENCH_DEFINE(BenchZeroB, OpDataBasic, zero_b, state.range(1))
+MPC_BENCH_DEFINE(BenchRandA, OpDataBasic, rand_a, Shape{state.range(1)})
+MPC_BENCH_DEFINE(BenchRandB, OpDataBasic, rand_b, Shape{state.range(1)})
 MPC_BENCH_DEFINE(BenchP2A, OpData1P, p2a, ps[0])
 MPC_BENCH_DEFINE(BenchA2P, OpData1A, a2p, as[0])
 MPC_BENCH_DEFINE(BenchMsbA2b, OpData1A, msb_a2b, as[0])
diff --git a/libspu/mpc/tools/complexity.cc b/libspu/mpc/tools/complexity.cc
index df28833d..cfd55bb8 100644
--- a/libspu/mpc/tools/complexity.cc
+++ b/libspu/mpc/tools/complexity.cc
@@ -23,10 +23,10 @@
 #include "llvm/Support/CommandLine.h"
 #include "spdlog/spdlog.h"
 
+#include "libspu/mpc/ab_api.h"
 #include "libspu/mpc/api.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/factory.h"
-#include "libspu/mpc/object.h"
 #include "libspu/mpc/utils/simulate.h"
 
 #include "libspu/mpc/tools/complexity.pb.h"
@@ -50,16 +50,18 @@ internal::SingleComplexityReport dumpComplexityReport(
 
   RuntimeConfig rt_conf;
   rt_conf.set_protocol(protocol);
+  rt_conf.set_field(FM64);
 
   utils::simulate(
       party_cnt, [&](const std::shared_ptr<yacl::link::Context>& lctx) -> void {
-        auto prot = Factory::CreateCompute(rt_conf, lctx);
+        SPUContext sctx(rt_conf, lctx);
+        Factory::RegisterProtocol(&sctx, lctx);
         if (lctx->Rank() != 0) {
           return;
         }
 
         for (auto name : kWhitelist) {
-          auto* kernel = prot->getKernel(name);
+          auto* kernel = sctx.getKernel(name);
           auto latency = kernel->latency();
           auto comm = kernel->comm();
 
diff --git a/libspu/mpc/utils/BUILD.bazel b/libspu/mpc/utils/BUILD.bazel
index 5594365d..017ef23b 100644
--- a/libspu/mpc/utils/BUILD.bazel
+++ b/libspu/mpc/utils/BUILD.bazel
@@ -51,23 +51,6 @@ spu_cc_test(
     ],
 )
 
-spu_cc_library(
-    name = "cexpr",
-    srcs = ["cexpr.cc"],
-    hdrs = ["cexpr.h"],
-    deps = [
-        "//libspu/core",
-    ],
-)
-
-spu_cc_test(
-    name = "cexpr_test",
-    srcs = ["cexpr_test.cc"],
-    deps = [
-        ":cexpr",
-    ],
-)
-
 spu_cc_library(
     name = "ring_ops",
     srcs = ["ring_ops.cc"],
@@ -120,3 +103,12 @@ spu_cc_test(
         ":linalg",
     ],
 )
+
+spu_cc_library(
+    name = "tiling_util",
+    hdrs = ["tiling_util.h"],
+    deps = [
+        "//libspu/core:context",
+        "//libspu/core:parallel_utils",
+    ],
+)
diff --git a/libspu/mpc/utils/tiling_util.h b/libspu/mpc/utils/tiling_util.h
new file mode 100644
index 00000000..78416115
--- /dev/null
+++ b/libspu/mpc/utils/tiling_util.h
@@ -0,0 +1,181 @@
+// Copyright 2023 Ant Group Co., Ltd.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except x 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 x 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 <future>
+
+#include "libspu/core/context.h"
+#include "libspu/core/parallel_utils.h"
+#include "libspu/core/prelude.h"
+
+namespace spu::mpc {
+
+// Inituition: some of the ops has complicated (compute, communication, compute,
+// communication ...) behaviour, when the input is large, CPU will wait for comm
+// interleavely. Pipeline is a software instruction level parallelism that use
+// tiling+concurrent to reduce the waiting time.
+//
+// Tiling+concurrent could be treated as the opposite of fusion+vectorization.
+
+namespace detail {
+
+template <typename Fn, typename... Args>
+ArrayRef tiled(Fn&& fn, SPUContext* ctx, const ArrayRef& x, Args&&... args) {
+  const int64_t kBlockSize = kMinTaskSize;
+  if (!ctx->config().experimental_enable_intra_op_par()  //
+      || !ctx->prot()->hasLowCostFork()                  //
+      || x.numel() <= kBlockSize                         //
+  ) {
+    return fn(ctx, x, std::forward<Args>(args)...);
+  }
+
+  std::vector<std::unique_ptr<SPUContext>> sub_ctxs;
+
+  const int64_t numBlocks =
+      x.numel() / kBlockSize + ((x.numel() % kBlockSize) != 0 ? 1 : 0);
+
+  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
+    sub_ctxs.push_back(ctx->fork());
+  }
+
+  std::vector<std::future<ArrayRef>> futures;
+  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
+    auto async_res = std::async(
+        [&](int64_t index) {
+          int64_t begin = index * kBlockSize;
+          int64_t end = std::min(begin + kBlockSize, x.numel());
+
+          return fn(sub_ctxs[index].get(), x.slice(begin, end), args...);
+        },
+        blk_idx);
+    futures.push_back(std::move(async_res));
+  }
+
+  std::vector<ArrayRef> out_slices;
+  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
+    out_slices.push_back(futures[blk_idx].get());
+  }
+
+  // Assume out.numel = x.numel
+  ArrayRef out(out_slices[0].eltype(), x.numel());
+  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
+    int64_t begin = blk_idx * kBlockSize;
+    int64_t end = std::min(begin + kBlockSize, x.numel());
+    std::memcpy(&out.at(begin), &out_slices[blk_idx].at(0),
+                (end - begin) * out.elsize());
+  }
+
+  return out;
+}
+
+template <typename Fn, typename... Args>
+ArrayRef tiled_2(Fn&& fn, SPUContext* ctx, const ArrayRef& x, const ArrayRef& y,
+                 Args&&... args) {
+  SPU_ENFORCE(x.numel() == y.numel());
+  const int64_t kBlockSize = kMinTaskSize;
+  if (!ctx->config().experimental_enable_intra_op_par()  //
+      || !ctx->prot()->hasLowCostFork()                  //
+      || x.numel() <= kBlockSize                         //
+  ) {
+    return fn(ctx, x, y, std::forward<Args>(args)...);
+  }
+
+  const int64_t numel = x.numel();
+
+  std::vector<std::unique_ptr<SPUContext>> sub_ctxs;
+
+  const int64_t numBlocks =
+      numel / kBlockSize + ((numel % kBlockSize) != 0 ? 1 : 0);
+
+  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
+    sub_ctxs.push_back(ctx->fork());
+  }
+
+  std::vector<std::future<ArrayRef>> futures;
+  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
+    auto async_res = std::async(
+        [&](int64_t index) {
+          int64_t begin = index * kBlockSize;
+          int64_t end = std::min(begin + kBlockSize, numel);
+
+          return fn(sub_ctxs[index].get(), x.slice(begin, end),
+                    y.slice(begin, end), args...);
+        },
+        blk_idx);
+    futures.push_back(std::move(async_res));
+  }
+
+  std::vector<ArrayRef> out_slices;
+  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
+    out_slices.push_back(futures[blk_idx].get());
+  }
+
+  // Assume out.numel = numel
+  ArrayRef out(out_slices[0].eltype(), numel);
+  for (int64_t blk_idx = 0; blk_idx < numBlocks; blk_idx++) {
+    int64_t begin = blk_idx * kBlockSize;
+    int64_t end = std::min(begin + kBlockSize, numel);
+    std::memcpy(&out.at(begin), &out_slices[blk_idx].at(0),
+                (end - begin) * out.elsize());
+  }
+
+  return out;
+}
+
+}  // namespace detail
+
+template <typename Fn, typename... Args>
+Value tiled(Fn&& fn, SPUContext* ctx, const Value& x, Args&&... args) {
+  // TODO: using nd-slice for tiling.
+  ArrayRef flat_x = flatten(x.data());
+
+  auto wrap_fn = [fn](SPUContext* wctx, const ArrayRef& wx,
+                      Args&&... wargs) -> ArrayRef {
+    auto res =
+        fn(wctx, WrapValue(wx, {wx.numel()}), std::forward<Args>(wargs)...);
+    auto [res_arr, res_shape, res_dtype] = UnwrapValue(res);
+    return res_arr;
+  };
+
+  auto flat_z =
+      detail::tiled(wrap_fn, ctx, flat_x, std::forward<Args>(args)...);
+  return WrapValue(flat_z, x.shape());
+}
+
+template <typename Fn, typename... Args>
+Value tiled(Fn&& fn, SPUContext* ctx, const Value& x, const Value& y,
+            Args&&... args) {
+  SPU_ENFORCE(x.shape() == y.shape());
+
+  // TODO: using nd-slice for tiling.
+  ArrayRef flat_x = flatten(x.data());
+  SPU_ENFORCE(flat_x.numel() == x.numel());
+  ArrayRef flat_y = flatten(y.data());
+  SPU_ENFORCE(flat_y.numel() == y.numel());
+
+  auto wrap_fn = [fn](SPUContext* wctx, const ArrayRef& wx, const ArrayRef& wy,
+                      Args&&... wargs) -> ArrayRef {
+    auto res = fn(wctx, WrapValue(wx, {wx.numel()}),
+                  WrapValue(wy, {wy.numel()}), std::forward<Args>(wargs)...);
+    auto [res_arr, res_shape, res_dtype] = UnwrapValue(res);
+    return res_arr;
+  };
+
+  // FIXME: if not rename to tile_2, it will select the first (less specialized)
+  // candidate.
+  auto flat_z = detail::tiled_2(wrap_fn, ctx, flat_x, flat_y,
+                                std::forward<Args>(args)...);
+  return WrapValue(flat_z, x.shape());
+}
+
+}  // namespace spu::mpc
diff --git a/libspu/spu.proto b/libspu/spu.proto
index a2b3511d..3d584257 100644
--- a/libspu/spu.proto
+++ b/libspu/spu.proto
@@ -303,10 +303,44 @@ message TTPBeaverConfig {
 //////////////////////////////////////////////////////////////////////////
 // Compiler relate definition
 //////////////////////////////////////////////////////////////////////////
+enum SourceIRType {
+  XLA = 0;
+  MLIR_HLO = 1;
+}
+
+message CompilationSource {
+  // Input IR type
+  SourceIRType ir_type = 1;
+
+  // IR
+  bytes ir_txt = 2;
+
+  // Input visibilities
+  repeated Visibility input_visibility = 3;
+}
+
+message CompilerOptions {
+  // Pretty print
+  bool enable_pretty_print = 1;
+  string pretty_print_dump_dir = 2;
+  
+  // Disable sqrt(x) + eps to sqrt(x+eps) rewrite
+  bool disable_sqrt_plus_epsilon_rewrite = 3;
+
+  // Disable x/sqrt(y) to x*rsqrt(y) rewrite
+  bool disable_div_sqrt_rewrite = 4;
+
+  // Disable reduce truncation optimization
+  bool disable_reduce_truncation_optimization = 5;
+
+  // Disable maxpooling optimization
+  bool disable_maxpooling_optimization = 6;
+
+  // Disallow mix type operations
+  bool disallow_mix_types_opts = 7;
 
-// Internal representation used by compiler
-message XlaMeta {
-  repeated Visibility inputs = 1;
+  // Disable SelectOp optimization
+  bool disable_select_optimization = 8;
 }
 
 // The executable format accepted by SPU runtime.
diff --git a/spu/api.py b/spu/api.py
index 4e7890eb..39b9f227 100644
--- a/spu/api.py
+++ b/spu/api.py
@@ -131,25 +131,19 @@ def reconstruct(self, str_shares: List[bytes]) -> 'np.ndarray':
 
 
 @cached(cache=LRUCache(maxsize=128))
-def _spu_compilation(ir_text: str, ir_type: str, json_meta: str):
-    pp_dir = os.getenv('SPU_IR_DUMP_DIR')
-    return libspu.compile(ir_text, ir_type, json_meta, pp_dir or "")
+def _spu_compilation(source: str, options_str: str):
+    return libspu.compile(source, options_str)
 
 
-def compile(ir_text: str, ir_type: str, vis: List[spu_pb2.Visibility]) -> str:
+def compile(source: spu_pb2.CompilationSource, copts: spu_pb2.CompilerOptions) -> str:
     """Compile from textual HLO/MHLO IR to SPU bytecode.
 
     Args:
-        ir_text (str): textual HLO/MHLO IR protobuf binary format.
-        ir_type (str): "hlo" or "mhlo".
-        vtype (spu_pb2.Visibility): Visbilities .
+        source (spu_pb2.CompilationSource): input to compiler.
+        copts (spu_pb2.CompilerOptions): compiler options.
 
     Returns:
         [spu_pb2.ValueProto]: output.
     """
-    from google.protobuf.json_format import MessageToJson
 
-    # todo: rename spu_pb2.XlaMeta to IrMeta?
-    return _spu_compilation(
-        ir_text, ir_type, MessageToJson(spu_pb2.XlaMeta(inputs=vis))
-    )
+    return _spu_compilation(source.SerializeToString(), copts.SerializeToString())
diff --git a/spu/libspu.cc b/spu/libspu.cc
index c8132e41..0810e876 100644
--- a/spu/libspu.cc
+++ b/spu/libspu.cc
@@ -24,13 +24,15 @@
 
 #include "libspu/compiler/common/compilation_context.h"
 #include "libspu/compiler/compile.h"
+#include "libspu/core/config.h"
+#include "libspu/core/context.h"
 #include "libspu/core/logging.h"
 #include "libspu/core/type_util.h"
+#include "libspu/core/value.h"
 #include "libspu/device/api.h"
 #include "libspu/device/io.h"
 #include "libspu/device/pphlo/pphlo_executor.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/value.h"
+#include "libspu/mpc/factory.h"
 #include "libspu/pir/pir.h"
 #include "libspu/psi/bucket_psi.h"
 #include "libspu/psi/core/ecdh_psi.h"
@@ -251,7 +253,7 @@ void BindLink(py::module& m) {
 
 // Wrap Runtime, it's workaround for protobuf pybind11/protoc conflict.
 class RuntimeWrapper {
-  std::unique_ptr<spu::HalContext> hctx_;
+  std::unique_ptr<spu::SPUContext> sctx_;
 
   // the golbals, could be used to cross session stuffs.
   spu::device::SymbolTable env_;
@@ -262,7 +264,11 @@ class RuntimeWrapper {
     spu::RuntimeConfig config;
     SPU_ENFORCE(config.ParseFromString(config_pb));
 
-    hctx_ = std::make_unique<spu::HalContext>(config, lctx);
+    // first, fill protobuf default value with implementation defined value.
+    populateRuntimecConfig(config);
+
+    sctx_ = std::make_unique<spu::SPUContext>(config, lctx);
+    mpc::Factory::RegisterProtocol(sctx_.get(), lctx);
   }
 
   void Run(const py::bytes& exec_pb) {
@@ -270,7 +276,7 @@ class RuntimeWrapper {
     SPU_ENFORCE(exec.ParseFromString(exec_pb));
 
     spu::device::pphlo::PPHloExecutor executor;
-    spu::device::execute(&executor, hctx_.get(), exec, &env_);
+    spu::device::execute(&executor, sctx_.get(), exec, &env_);
   }
 
   void SetVar(const std::string& name, const py::bytes& value) {
@@ -591,24 +597,18 @@ PYBIND11_MODULE(libspu, m) {
   // bind compiler.
   m.def(
       "compile",
-      [](const py::bytes& ir_text, const std::string& ir_type,
-         const std::string& input_visiblity_map, const std::string& dump_path) {
+      [](const py::bytes& source, const std::string& copts) {
         py::scoped_ostream_redirect stream(
             std::cout,                                 // std::ostream&
             py::module_::import("sys").attr("stdout")  // Python output
         );
 
         spu::compiler::CompilationContext ctx;
-        ctx.setInputVisibilityString(input_visiblity_map);
-
-        if (!dump_path.empty()) {
-          ctx.enablePrettyPrintWithDir(dump_path);
-        }
+        ctx.setCompilerOptions(copts);
 
-        return py::bytes(spu::compiler::compile(&ctx, ir_text, ir_type));
+        return py::bytes(spu::compiler::compile(&ctx, source));
       },
-      "spu compile.", py::arg("ir_text"), py::arg("ir_type"),
-      py::arg("vis_map"), py::arg("dump_path"));
+      "spu compile.", py::arg("source"), py::arg("copts"));
 
   // bind spu libs.
   py::module link_m = m.def_submodule("link");
diff --git a/spu/utils/distributed.py b/spu/utils/distributed.py
index 9c051939..811d4c0e 100644
--- a/spu/utils/distributed.py
+++ b/spu/utils/distributed.py
@@ -644,16 +644,17 @@ def __del__(self):
     class JaxFunction(Device.Function):
         device: SPU
 
-        def __init__(self, device: SPU, pyfunc: Callable, static_argnums):
+        def __init__(self, device: SPU, pyfunc: Callable, static_argnums, copts):
             super().__init__(device, pyfunc)
             self.static_argnums = static_argnums
+            self.copts = copts
 
         def __call__(self, *args, **kwargs):
             args, kwargs = self.device._place_arguments(*args, **kwargs)
 
             # now, all object are either PyObject or SPU.DeviceObject
             executable, args_flat, out_tree = self._compile_jax_func(
-                self.pyfunc, self.static_argnums, *args, **kwargs
+                self.pyfunc, self.static_argnums, self.copts, *args, **kwargs
             )
 
             def get_share_ref(idx, obj):
@@ -688,11 +689,11 @@ def get_share_ref(idx, obj):
         def dump_pphlo(self, *args, **kwargs):
             args, kwargs = self.device._place_arguments(*args, **kwargs)
             executable, *_ = self._compile_jax_func(
-                self.pyfunc, self.static_argnums, *args, **kwargs
+                self.pyfunc, self.static_argnums, self.copts, *args, **kwargs
             )
             return executable.code.decode('utf-8')
 
-        def _compile_jax_func(self, fn, static_argnums, *args, **kwargs):
+        def _compile_jax_func(self, fn, static_argnums, copts, *args, **kwargs):
             def mock_parameters(obj: Union[SPU.Object, np.ndarray]):
                 if isinstance(obj, SPU.Object):
                     return np.zeros(shape=obj.shape, dtype=obj.dtype)
@@ -716,6 +717,7 @@ def mock_parameters(obj: Union[SPU.Object, np.ndarray]):
                 else spu_pb2.Visibility.VIS_PUBLIC
                 for arg in args_flat
             ]
+
             in_names = [f'{id(fn_name)}-in{idx}' for idx in range(len(args_flat))]
 
             def outputNameGen(out_flat: List):
@@ -730,6 +732,7 @@ def outputNameGen(out_flat: List):
                 in_vis,
                 outputNameGen,
                 static_argnums=static_argnums,
+                copts=copts,
             )
 
             _, output_tree = jax.tree_util.tree_flatten(output)
@@ -738,15 +741,18 @@ def outputNameGen(out_flat: List):
     class TensorFlowFunction(Device.Function):
         device: SPU
 
-        def __init__(self, device: Device, pyfunc: Callable):
+        def __init__(
+            self, device: Device, pyfunc: Callable, copts: spu_pb2.CompilerOptions
+        ):
             super().__init__(device, pyfunc)
+            self.copts = copts
 
         def __call__(self, *args, **kwargs):
             args, kwargs = self.device._place_arguments(*args, **kwargs)
 
             # now, all object are either PyObject or SPU.DeviceObject
             pphlo_ir, args_flat, structured_outputs = self._compile_tf_func(
-                self.pyfunc, *args, **kwargs
+                self.pyfunc, self.copts, *args, **kwargs
             )
 
             def get_share_ref(idx, obj):
@@ -784,7 +790,7 @@ def get_share_ref(idx, obj):
                 structured_outputs, ret_flat, expand_composites=True
             )
 
-        def _compile_tf_func(self, fn, *args, **kwargs):
+        def _compile_tf_func(self, fn, copts, *args, **kwargs):
             def mock_parameters(obj: Union[SPU.Object, np.ndarray]):
                 if isinstance(obj, SPU.Object):
                     return np.zeros(shape=obj.shape, dtype=obj.dtype)
@@ -806,6 +812,7 @@ def mock_parameters(obj: Union[SPU.Object, np.ndarray]):
                 else spu_pb2.Visibility.VIS_PUBLIC
                 for arg in args_flat
             ]
+
             in_names = [f'{id(fn_name)}-in{idx}' for idx in range(len(args_flat))]
 
             def outputNameGen(out_flat: List):
@@ -819,6 +826,7 @@ def outputNameGen(out_flat: List):
                 in_names,
                 in_vis,
                 outputNameGen,
+                copts=copts,
             )
             return executable, args_flat, output_tree
 
@@ -949,11 +957,16 @@ def details(self) -> str:
         hosts = [nc.addr for nc in self.node_clients]
         return f"name: {self.name}\nhosted by: {hosts}\ninternal addrs: {self.internal_addrs}\n{self.runtime_config}"
 
-    def compile(self, fn: Callable, static_argnums=()) -> Callable:
+    def compile(
+        self,
+        fn: Callable,
+        static_argnums=(),
+        copts=spu_pb2.CompilerOptions(),
+    ) -> Callable:
         if _FRAMEWORK == Framework.EXP_TF:
-            return SPU.TensorFlowFunction(self, fn)
+            return SPU.TensorFlowFunction(self, fn, copts)
         elif _FRAMEWORK == Framework.JAX:
-            return SPU.JaxFunction(self, fn, static_argnums)
+            return SPU.JaxFunction(self, fn, static_argnums, copts)
         elif _FRAMEWORK == Framework.EXP_TORCH:
             return SPU.TorchFunction(self, fn)
         else:
diff --git a/spu/utils/frontend.py b/spu/utils/frontend.py
index 209d80af..d810f292 100644
--- a/spu/utils/frontend.py
+++ b/spu/utils/frontend.py
@@ -102,6 +102,7 @@ def compile(
     input_vis: List,
     outputNameGen: Callable,
     static_argnums=(),
+    copts=spu_pb2.CompilerOptions(),
 ):
     if kind == Kind.JAX:
         import jax
@@ -167,13 +168,16 @@ def compile(
     else:
         raise NameError(f"Unknown frontend type {kind}")
 
+    source = spu_pb2.CompilationSource()
+    source.ir_txt = ir_text
+    source.input_visibility.extend(input_vis)
     if kind in [Kind.JAX, Kind.Tensorflow]:
-        ir_type = "hlo"
+        source.ir_type = spu_pb2.SourceIRType.XLA
         name = fn.func.__name__ if isinstance(fn, functools.partial) else fn.__name__
     elif kind == Kind.Torch:
-        ir_type = "mhlo"
+        source.ir_type = spu_pb2.SourceIRType.MLIR_HLO
         name = repr(fn)
-    mlir = spu_api.compile(ir_text, ir_type, input_vis)
+    mlir = spu_api.compile(source, copts)
     executable = spu_pb2.ExecutableProto(
         name=name,
         input_names=input_names,
diff --git a/spu/utils/simulation.py b/spu/utils/simulation.py
index d2ee3f6a..ac71a003 100644
--- a/spu/utils/simulation.py
+++ b/spu/utils/simulation.py
@@ -120,7 +120,7 @@ def sim_jax(
     sim: Simulator,
     fun: Callable,
     static_argnums=(),
-    inline: bool = False,
+    copts=spu_pb2.CompilerOptions(),
 ):
     """
     Decorates a jax numpy fn that simulated on SPU.
@@ -144,7 +144,6 @@ def wrapper(*args, **kwargs):
         args_flat, _ = jax.tree_util.tree_flatten((dyn_args, kwargs))
 
         in_names = [f'in{idx}' for idx in range(len(args_flat))]
-        in_vis = [spu_pb2.Visibility.VIS_SECRET] * len(args_flat)
 
         def outputNameGen(out_flat):
             return [f'out{idx}' for idx in range(len(out_flat))]
@@ -155,9 +154,10 @@ def outputNameGen(out_flat):
             args,
             kwargs,
             in_names,
-            in_vis,
+            [spu_pb2.Visibility.VIS_SECRET] * len(args_flat),
             outputNameGen,
             static_argnums=static_argnums,
+            copts=copts,
         )
 
         wrapper.pphlo = executable.code.decode("utf-8")
diff --git a/spu/version.py b/spu/version.py
index 0a6a2730..49ab6fed 100644
--- a/spu/version.py
+++ b/spu/version.py
@@ -13,4 +13,4 @@
 # limitations under the License.
 
 
-__version__ = "0.3.3b3"
+__version__ = "0.4.0b0"


Test patch:
diff --git a/libspu/compiler/tests/BUILD.bazel b/libspu/compiler/tests/BUILD.bazel
index fdee767d..1649fc05 100644
--- a/libspu/compiler/tests/BUILD.bazel
+++ b/libspu/compiler/tests/BUILD.bazel
@@ -13,6 +13,7 @@
 # limitations under the License.
 
 load("//bazel:glob_lit_test.bzl", "glob_lit_tests")
+load("//bazel:spu.bzl", "spu_cc_test")
 
 glob_lit_tests(
     data = [":test_utilities"],
@@ -30,3 +31,12 @@ filegroup(
         "@llvm-project//llvm:FileCheck",
     ],
 )
+
+spu_cc_test(
+    name = "enum_conversion_test",
+    srcs = ["enum_conversion_test.cc"],
+    deps = [
+        "//libspu:spu_cc_proto",
+        "//libspu/dialect:pphlo_dialect",
+    ],
+)
diff --git a/libspu/compiler/tests/enum_conversion_test.cc b/libspu/compiler/tests/enum_conversion_test.cc
new file mode 100644
index 00000000..aa9018ee
--- /dev/null
+++ b/libspu/compiler/tests/enum_conversion_test.cc
@@ -0,0 +1,65 @@
+// Copyright 2023 Ant Group Co., Ltd.
+//
+// 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 "gtest/gtest.h"
+
+#include "libspu/dialect/pphlo_base_enums.h"
+
+#include "libspu/spu.pb.h"
+
+namespace spu {
+
+// Compiler options relies on pb and td has the same string representation of
+// fields, use a round trip test to lock this down
+TEST(EnumConversion, Public) {
+#define CHECK(T)                                                               \
+  auto v = Visibility::T;                                                      \
+  auto mlir_v =                                                                \
+      mlir::pphlo::symbolizeEnum<mlir::pphlo::Visibility>(Visibility_Name(v)); \
+  EXPECT_EQ(mlir_v, mlir::pphlo::Visibility::T);
+
+  {CHECK(VIS_PUBLIC)} { CHECK(VIS_SECRET) }
+
+#undef CHECK
+}
+
+TEST(EnumConversion, ProtoKinds) {
+  Visibility v;
+  // This is basically a compile time check...that there is no new field
+  switch (v) {
+  case Visibility::VIS_SECRET:
+    break;
+  case Visibility::VIS_PUBLIC:
+    break;
+  case Visibility::VIS_INVALID:
+    break;
+  case Visibility::Visibility_INT_MAX_SENTINEL_DO_NOT_USE_:
+    break;
+  case Visibility::Visibility_INT_MIN_SENTINEL_DO_NOT_USE_:
+    break;
+  }
+}
+
+TEST(EnumConversion, TdKinds) {
+  mlir::pphlo::Visibility v;
+  // This is basically a compile time check...that there is no new field
+  switch (v) {
+  case mlir::pphlo::Visibility::VIS_SECRET:
+    break;
+  case mlir::pphlo::Visibility::VIS_PUBLIC:
+    break;
+  }
+}
+
+} // namespace spu
diff --git a/libspu/compiler/tests/hlo_to_pphlo_comparison.mlir b/libspu/compiler/tests/hlo_to_pphlo_comparison.mlir
index ee7380ff..d4cc6044 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_comparison.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_comparison.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo --split-input-file %s | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_PUBLIC,VIS_PUBLIC --split-input-file %s | FileCheck %s
 
 func.func @main(%arg0: tensor<2x2xi32>, %arg1: tensor<2x2xi32>) -> (tensor<2x2xi1>) {
     // CHECK: "pphlo.equal"(%arg0, %arg1) : (tensor<2x2x!pphlo.pub<i32>>, tensor<2x2x!pphlo.pub<i32>>) -> tensor<2x2x!pphlo.pub<i1>>
diff --git a/libspu/compiler/tests/hlo_to_pphlo_conditional.mlir b/libspu/compiler/tests/hlo_to_pphlo_conditional.mlir
index 2e4b83b4..ee27150e 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_conditional.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_conditional.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt -hlo-legalize-to-pphlo %s --split-input-file  | FileCheck %s
+// RUN: mlir-pphlo-opt -hlo-legalize-to-pphlo=input_vis_list=VIS_PUBLIC %s --split-input-file  | FileCheck %s
 
 func.func @main(%arg0: tensor<f32>) -> tensor<f32> {
   //CHECK: %2 = "pphlo.if"(%1) ({
diff --git a/libspu/compiler/tests/hlo_to_pphlo_conditional_mixed.mlir b/libspu/compiler/tests/hlo_to_pphlo_conditional_mixed.mlir
index 257261c8..6c1c32f7 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_conditional_mixed.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_conditional_mixed.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt -hlo-legalize-to-pphlo='io-visibility-json={"inputs":["VIS_PUBLIC","VIS_SECRET"]}' %s --split-input-file  | FileCheck %s
+// RUN: mlir-pphlo-opt -hlo-legalize-to-pphlo=input_vis_list=VIS_PUBLIC,VIS_SECRET %s --split-input-file  | FileCheck %s
 
 func.func @main(%arg0: tensor<f32>, %arg1: tensor<f32>) -> tensor<f32> {
   //CHECK: %2 = "pphlo.if"(%1) ({
diff --git a/libspu/compiler/tests/hlo_to_pphlo_conditional_secret_cond.mlir b/libspu/compiler/tests/hlo_to_pphlo_conditional_secret_cond.mlir
index 3b36c744..7a9b7f22 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_conditional_secret_cond.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_conditional_secret_cond.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt -hlo-legalize-to-pphlo='io-visibility-json={"inputs":["VIS_SECRET"]}' %s --split-input-file  | FileCheck %s
+// RUN: mlir-pphlo-opt -hlo-legalize-to-pphlo=input_vis_list=VIS_SECRET %s --split-input-file  | FileCheck %s
 
 func.func @main(%arg0: tensor<f32>) -> tensor<f32> {
   //CHECK: %2 = "pphlo.if"(%1) ({
diff --git a/libspu/compiler/tests/hlo_to_pphlo_ops_binary.mlir b/libspu/compiler/tests/hlo_to_pphlo_ops_binary.mlir
index c0586510..81c14538 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_ops_binary.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_ops_binary.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo --split-input-file %s | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_PUBLIC,VIS_PUBLIC --split-input-file %s | FileCheck %s
 
 func.func @main(%arg0: tensor<2x2xi32>, %arg1: tensor<2x2xi32>) -> (tensor<2x2xi32>) {
     // CHECK: %0 = "pphlo.subtract"(%arg0, %arg1) : (tensor<2x2x!pphlo.pub<i32>>, tensor<2x2x!pphlo.pub<i32>>) -> tensor<2x2x!pphlo.pub<i32>>
diff --git a/libspu/compiler/tests/hlo_to_pphlo_ops_other.mlir b/libspu/compiler/tests/hlo_to_pphlo_ops_other.mlir
index ff8e4c55..878e6448 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_ops_other.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_ops_other.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo --split-input-file %s | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_PUBLIC,VIS_PUBLIC,VIS_PUBLIC,VIS_PUBLIC,VIS_PUBLIC --split-input-file %s | FileCheck %s
 
 func.func @main(%arg0: tensor<16xf32>,%arg1: tensor<1024x1xi1>, %arg2: tensor<1024x1xf32>, %arg3: tensor<1024x1xf32>, %arg4: tensor<3x4xi32>) -> (tensor<1024x16xf32>) {
     // CHECK: %0 = "pphlo.broadcast"(%arg0) {broadcast_dimensions = dense<1> : tensor<1xi64>} : (tensor<16x!pphlo.pub<f32>>) -> tensor<1024x16x!pphlo.pub<f32>>
diff --git a/libspu/compiler/tests/hlo_to_pphlo_ops_unary.mlir b/libspu/compiler/tests/hlo_to_pphlo_ops_unary.mlir
index 6d1b3400..d6215485 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_ops_unary.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_ops_unary.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo --split-input-file %s | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_PUBLIC --split-input-file %s | FileCheck %s
 
 func.func @main(%arg0: tensor<2x2xi32>) -> (tensor<2x2xf32>) {
     // CHECK:  %0 = "pphlo.convert"(%arg0) : (tensor<2x2x!pphlo.pub<i32>>) -> tensor<2x2x!pphlo.pub<f32>>
diff --git a/libspu/compiler/tests/hlo_to_pphlo_select_and_scatter.mlir b/libspu/compiler/tests/hlo_to_pphlo_select_and_scatter.mlir
index 45bb7de9..e60e4bde 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_select_and_scatter.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_select_and_scatter.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo='io-visibility-json={"inputs":["VIS_SECRET"]}' --split-input-file %s | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_SECRET,VIS_PUBLIC,VIS_PUBLIC --split-input-file %s | FileCheck %s
 
 func.func @main(%arg0: tensor<128x5x5x32xf32>, %arg1: tensor<128x4x4x32xf32>, %arg2: tensor<f32>) -> tensor<128x5x5x32xf32>   {
   // CHECK: %1 = "pphlo.select_and_scatter"(%arg0, %arg1, %0) ({
diff --git a/libspu/compiler/tests/hlo_to_pphlo_sort.mlir b/libspu/compiler/tests/hlo_to_pphlo_sort.mlir
index 2e0654f1..6b6e9acf 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_sort.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_sort.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo='io-visibility-json={"inputs":["VIS_SECRET"]}' --split-input-file %s | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_SECRET --split-input-file %s | FileCheck %s
 
 func.func @main(%arg0: tensor<20xi32>) -> (tensor<20xi32>) {
     %0 = stablehlo.iota dim = 0 : tensor<20xi32>
diff --git a/libspu/compiler/tests/hlo_to_pphlo_while.mlir b/libspu/compiler/tests/hlo_to_pphlo_while.mlir
index e74b384d..1e31018f 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_while.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_while.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt -hlo-legalize-to-pphlo %s --split-input-file  | FileCheck %s
+// RUN: mlir-pphlo-opt -hlo-legalize-to-pphlo=input_vis_list=VIS_PUBLIC %s --split-input-file  | FileCheck %s
 
 func.func @main(%arg0: tensor<i64>) -> tensor<i64> {
   //CHECK: %0 = "pphlo.while"(%arg0) ({
diff --git a/libspu/compiler/tests/pphlo_type_inference_all_public.mlir b/libspu/compiler/tests/pphlo_type_inference_all_public.mlir
index cc33aa90..4b4b2bca 100644
--- a/libspu/compiler/tests/pphlo_type_inference_all_public.mlir
+++ b/libspu/compiler/tests/pphlo_type_inference_all_public.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo='io-visibility-json={"inputs":["VIS_PUBLIC"]}' %s --split-input-file  | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_PUBLIC,VIS_PUBLIC %s --split-input-file  | FileCheck %s
 
 func.func @main(%arg0: tensor<2x2xf32>, %arg1: tensor<2x2xf32>) -> (tensor<2x2xui8>) {
     // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui8>} : () -> tensor<2x2x!pphlo.pub<ui8>> 
diff --git a/libspu/compiler/tests/pphlo_type_inference_all_secret.mlir b/libspu/compiler/tests/pphlo_type_inference_all_secret.mlir
index 3a74cf17..bede2bba 100644
--- a/libspu/compiler/tests/pphlo_type_inference_all_secret.mlir
+++ b/libspu/compiler/tests/pphlo_type_inference_all_secret.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo='io-visibility-json={"inputs":["VIS_SECRET","VIS_SECRET"]}' %s --split-input-file  | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_SECRET,VIS_SECRET %s --split-input-file  | FileCheck %s
 
 // CHECK: func @main(%arg0: tensor<2x2x!pphlo.sec<f32>>, %arg1: tensor<2x2x!pphlo.sec<f32>>) -> tensor<2x2x!pphlo.sec<f32>> {
 func.func @main(%arg0: tensor<2x2xf32>, %arg1: tensor<2x2xf32>) -> (tensor<2x2xf32>) {
diff --git a/libspu/compiler/tests/pphlo_type_inference_identity.mlir b/libspu/compiler/tests/pphlo_type_inference_identity.mlir
index 7745c375..5ebee12f 100644
--- a/libspu/compiler/tests/pphlo_type_inference_identity.mlir
+++ b/libspu/compiler/tests/pphlo_type_inference_identity.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo='io-visibility-json={"inputs":["VIS_SECRET","VIS_SECRET"]}' %s --split-input-file  | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_SECRET,VIS_SECRET %s --split-input-file  | FileCheck %s
 
 func.func @main(%arg0: tensor<10xf64>, %arg1: tensor<10xf64>) -> (tensor<10xf64>, tensor<10xf64>) {
     // CHECK: return %arg0, %arg1 : tensor<10x!pphlo.sec<f64>>, tensor<10x!pphlo.sec<f64>>
diff --git a/libspu/compiler/tests/pphlo_type_inference_reduce.mlir b/libspu/compiler/tests/pphlo_type_inference_reduce.mlir
index b509382e..d50af752 100644
--- a/libspu/compiler/tests/pphlo_type_inference_reduce.mlir
+++ b/libspu/compiler/tests/pphlo_type_inference_reduce.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo='io-visibility-json={"inputs":["VIS_SECRET"]}' %s --split-input-file  | FileCheck %s
+// RUN: mlir-pphlo-opt --hlo-legalize-to-pphlo=input_vis_list=VIS_SECRET %s --split-input-file  | FileCheck %s
 
 // CHECK: func @main(%arg0: tensor<1024x1x!pphlo.sec<f32>>) -> tensor<1024x!pphlo.sec<f32>> { 
 func.func @main(%arg1: tensor<1024x1xf32>) -> (tensor<1024xf32>) {
diff --git a/libspu/mpc/utils/cexpr_test.cc b/libspu/core/cexpr_test.cc
similarity index 94%
rename from libspu/mpc/utils/cexpr_test.cc
rename to libspu/core/cexpr_test.cc
index 1c9dbd45..546ef8ef 100644
--- a/libspu/mpc/utils/cexpr_test.cc
+++ b/libspu/core/cexpr_test.cc
@@ -12,13 +12,13 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "libspu/mpc/utils/cexpr.h"
+#include "libspu/core/cexpr.h"
 
 #include <cmath>
 
 #include "gtest/gtest.h"
 
-namespace spu::mpc::cexpr {
+namespace spu::ce {
 
 TEST(CExprTest, K) {
   EXPECT_EQ(K()->expr(), "K");
@@ -56,4 +56,4 @@ TEST(CExprTest, Samples) {
   EXPECT_EQ(c->eval({{"K", 32}, {"N", 2}}), 2 * (std::log2(32) - 1) + 3 * 2);
 }
 
-}  // namespace spu::mpc::cexpr
+}  // namespace spu::ce
diff --git a/libspu/kernel/value_test.cc b/libspu/core/value_test.cc
similarity index 98%
rename from libspu/kernel/value_test.cc
rename to libspu/core/value_test.cc
index 6602aaf1..57169d57 100644
--- a/libspu/kernel/value_test.cc
+++ b/libspu/core/value_test.cc
@@ -12,7 +12,7 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "libspu/kernel/value.h"
+#include "libspu/core/value.h"
 
 #include "gtest/gtest.h"
 
diff --git a/libspu/device/io_test.cc b/libspu/device/io_test.cc
index a0deda49..bf146992 100644
--- a/libspu/device/io_test.cc
+++ b/libspu/device/io_test.cc
@@ -105,8 +105,8 @@ TEST_P(ColocatedIoTest, Works) {
   hconf.set_field(std::get<2>(GetParam()));
 
   mpc::utils::simulate(kWorldSize, [&](auto lctx) {
-    HalContext hctx(hconf, lctx);
-    ColocatedIo cio(&hctx);
+    SPUContext sctx(hconf, lctx);
+    ColocatedIo cio(&sctx);
 
     // WHEN
     if (lctx->Rank() == 0) {
diff --git a/libspu/device/pphlo/pphlo_executor_test.cc b/libspu/device/pphlo/pphlo_executor_test.cc
index a78d07dd..3206c01c 100644
--- a/libspu/device/pphlo/pphlo_executor_test.cc
+++ b/libspu/device/pphlo/pphlo_executor_test.cc
@@ -30,6 +30,7 @@
 #include "libspu/device/api.h"
 #include "libspu/device/symbol_table.h"
 #include "libspu/device/test_utils.h"
+#include "libspu/kernel/test_util.h"
 #include "libspu/mpc/ref2k/ref2k.h"
 #include "libspu/mpc/utils/simulate.h"
 
@@ -56,12 +57,16 @@ class Runner {
   }
 
   static std::string compileMHlo(const std::string &mhlo,
-                                 const std::string &vis = {}) {
-    compiler::CompilationContext ctx;
-    if (!vis.empty()) {
-      ctx.setInputVisibilityString(vis);
+                                 const std::vector<spu::Visibility> &vis) {
+    CompilationSource source;
+    source.set_ir_type(SourceIRType::MLIR_HLO);
+    source.set_ir_txt(mhlo);
+    for (const auto v : vis) {
+      source.add_input_visibility(v);
     }
-    return compiler::compile(&ctx, mhlo, "mhlo");
+
+    compiler::CompilationContext ctx;
+    return compiler::compile(&ctx, source.SerializeAsString());
   }
 
   void run(const std::string &mlir, size_t num_output = 1) {
@@ -76,10 +81,10 @@ class Runner {
           if (lctx->Rank() == 0) {
             // conf.set_enable_action_trace(true);
           }
-          HalContext hctx(conf, lctx);
+          SPUContext sctx = kernel::test::makeSPUContext(conf, lctx);
           auto *env = io_->GetSymbolTable(lctx->Rank());
           pphlo::PPHloExecutor executor;
-          execute(&executor, &hctx, executable_, env);
+          execute(&executor, &sctx, executable_, env);
         });
   }
 
@@ -809,7 +814,7 @@ void testGatherImpl(size_t world_size, FieldType field, ProtocolKind protocol,
     // Start indices
     r.addInput(indices);
 
-    auto compiled = r.compileMHlo(mhlo);
+    auto compiled = r.compileMHlo(mhlo, {VIS_PUBLIC, VIS_PUBLIC});
 
     EXPECT_THAT(compiled, testing::HasSubstr("pphlo.gather"));
 
@@ -826,8 +831,7 @@ void testGatherImpl(size_t world_size, FieldType field, ProtocolKind protocol,
     // Start indices
     r.addInput(indices, VIS_SECRET);
 
-    auto compiled =
-        r.compileMHlo(mhlo, R"({"inputs":["VIS_PUBLIC","VIS_SECRET"]})");
+    auto compiled = r.compileMHlo(mhlo, {VIS_PUBLIC, VIS_SECRET});
 
     EXPECT_THAT(compiled, testing::Not(testing::HasSubstr("pphlo.gather")));
 
@@ -944,7 +948,8 @@ func.func @main(%arg0: tensor<1x1x4x4xf32>, %arg1: tensor<1x1x2x2xf32>) -> (tens
               feature_group_count = 1 : i64
             } : (tensor<1x1x4x4xf32>, tensor<1x1x2x2xf32>) -> tensor<1x1x4x4xf32>
     return %0 : tensor<1x1x4x4xf32>
-})");
+})",
+                                             {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -982,7 +987,8 @@ func.func @main(%arg0: tensor<2x3x1x4xf32>, %arg1:tensor<1x3x2x3xf32>) -> (tenso
             feature_group_count = 1 : i64
           } : (tensor<2x3x1x4xf32>,tensor<1x3x2x3xf32>) -> tensor<1x1x1x2xf32>
     return %0 : tensor<1x1x1x2xf32>
-})");
+})",
+                                             {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -1019,7 +1025,8 @@ func.func @main(%arg0: tensor<1x1x4x4xf32>, %arg1: tensor<1x1x2x2xf32>) -> (tens
             feature_group_count = 1 : i64
           } : (tensor<1x1x4x4xf32>, tensor<1x1x2x2xf32>) -> tensor<1x1x7x7xf32>
     return %0 : tensor<1x1x7x7xf32>
-})");
+})",
+                                             {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -1062,7 +1069,8 @@ func.func @main(%arg0: tensor<1x1x4x4xf32>, %arg1: tensor<1x1x2x2xf32>) -> (tens
             feature_group_count = 1 : i64
           } : (tensor<1x1x4x4xf32>, tensor<1x1x2x2xf32>) -> tensor<1x1x8x8xf32>
     return %0 : tensor<1x1x8x8xf32>
-})");
+})",
+                                             {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -1107,7 +1115,8 @@ func.func @main(%arg0: tensor<1x1x4x6xf32>, %arg1: tensor<1x1x2x3xf32>) -> (tens
             feature_group_count = 1 : i64
           } : (tensor<1x1x4x6xf32>, tensor<1x1x2x3xf32>) -> tensor<1x1x2x2xf32>
     return %0 : tensor<1x1x2x2xf32>
-})");
+})",
+                                             {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -1964,7 +1973,8 @@ func.func @main(%arg0: tensor<4x6xi32>, %arg1: tensor<2x2xi32>) -> (tensor<2x2xi
       "mhlo.return"(%3) : (tensor<i32>) -> ()
     }) {padding = dense<0> : tensor<2x2xi64>, window_dimensions = dense<[2,3]> : tensor<2xi64>, window_strides = dense<[2,3]> : tensor<2xi64>} : (tensor<4x6xi32>, tensor<2x2xi32>, tensor<i32>) -> tensor<4x6xi32>
     return %1, %2 : tensor<2x2xi32>, tensor<4x6xi32>
-})");
+})",
+                                             {VIS_PUBLIC, VIS_PUBLIC});
 
   EXPECT_THAT(ir, testing::HasSubstr("pphlo.maxpool_scatter"));
 
@@ -2203,7 +2213,7 @@ func.func @main(%arg0: tensor<20xi32>) -> (tensor<20xi32>, tensor<20xi32>) {
     }) {dimension = 0 : i64, is_stable = true} : (tensor<20xi32>, tensor<20xi32>) -> (tensor<20xi32>, tensor<20xi32>)
     return %1#0, %1#1: tensor<20xi32>, tensor<20xi32>
 })",
-            R"({"inputs":["VIS_SECRET"]})"),
+            {VIS_SECRET}),
         2);
   r.verifyOutput(expected_ret0.data(), 0);
   r.verifyOutput(expected_ret1.data(), 1);
diff --git a/libspu/device/pphlo/pphlo_verifier_test.cc b/libspu/device/pphlo/pphlo_verifier_test.cc
index 33335306..90c73d14 100644
--- a/libspu/device/pphlo/pphlo_verifier_test.cc
+++ b/libspu/device/pphlo/pphlo_verifier_test.cc
@@ -21,6 +21,7 @@
 #include "xtensor/xarray.hpp"
 
 #include "libspu/device/test_utils.h"
+#include "libspu/kernel/test_util.h"
 #include "libspu/mpc/utils/simulate.h"
 
 namespace spu::device::pphlo {
@@ -49,8 +50,8 @@ void runner(const OpFcn &f, absl::Span<const xt::xarray<InT>> inputs,
 
   ::spu::mpc::utils::simulate(
       2, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-        HalContext hctx(conf, lctx);
-        PPHloVerifier verifier(&hctx);
+        SPUContext sctx = kernel::test::makeSPUContext(conf, lctx);
+        PPHloVerifier verifier(&sctx);
         verifier.setMismatchHandler(failed);
         auto *table = io_->GetSymbolTable(lctx->Rank());
 
@@ -267,8 +268,8 @@ TEST(Verify, Select) {
 
   ::spu::mpc::utils::simulate(
       2, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-        HalContext hctx(conf, lctx);
-        PPHloVerifier verifier(&hctx);
+        SPUContext sctx = kernel::test::makeSPUContext(conf, lctx);
+        PPHloVerifier verifier(&sctx);
         verifier.setMismatchHandler(failed);
         auto *table = io_->GetSymbolTable(lctx->Rank());
 
diff --git a/libspu/kernel/hal/array_element_wise_test.cc b/libspu/kernel/hal/array_element_wise_test.cc
index 09c7ede1..5a80328b 100644
--- a/libspu/kernel/hal/array_element_wise_test.cc
+++ b/libspu/kernel/hal/array_element_wise_test.cc
@@ -17,7 +17,7 @@
 #include "xtensor/xmath.hpp"
 
 #include "libspu/kernel/hal/polymorphic.h"
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
 
diff --git a/libspu/kernel/hal/constants_test.cc b/libspu/kernel/hal/constants_test.cc
index 57d6a460..7fd3ef67 100644
--- a/libspu/kernel/hal/constants_test.cc
+++ b/libspu/kernel/hal/constants_test.cc
@@ -17,12 +17,12 @@
 #include "gmock/gmock.h"
 #include "gtest/gtest.h"
 
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
 
 TEST(ConstantsTest, Scalar) {
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   Value i = constant(&ctx, 0, DT_I32);
   EXPECT_TRUE(i.shape().empty());
@@ -40,7 +40,7 @@ TEST(ConstantsTest, Scalar) {
 }
 
 TEST(ConstantsTest, Tensor) {
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> raw = {1.0F};
   Value x = constant(&ctx, raw, DT_FXP);
@@ -52,7 +52,7 @@ TEST(ConstantsTest, Tensor) {
 }
 
 TEST(ConstantsTest, TensorBroadcast) {
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> raw = {
       {1.0, 2.0},
@@ -67,7 +67,7 @@ TEST(ConstantsTest, TensorBroadcast) {
 }
 
 TEST(ConstantsTest, Initializer) {
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   // FIXME: the dtype is determined by the C++ literal type.
   // EXPECT_EQ(constant(&ctx, 0, DT_I1).dtype(), DT_I1);  // FIXME
diff --git a/libspu/kernel/hal/fxp_approx_test.cc b/libspu/kernel/hal/fxp_approx_test.cc
index 7798083e..64b5f0ae 100644
--- a/libspu/kernel/hal/fxp_approx_test.cc
+++ b/libspu/kernel/hal/fxp_approx_test.cc
@@ -18,12 +18,12 @@
 #include "xtensor/xio.hpp"
 
 #include "libspu/kernel/hal/constants.h"
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
 
 TEST(FxpTest, ExponentialPublic) {
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {
       -500.0, -100.0, -16.7, -14.3, -12.5, -11.0, -9.9, -6.7,
@@ -42,7 +42,7 @@ TEST(FxpTest, ExponentialPublic) {
 }
 
 TEST(FxpTest, ExponentialTaylorSeries) {
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   // GIVEN
   xt::xarray<float> x = {
@@ -63,7 +63,7 @@ TEST(FxpTest, ExponentialTaylorSeries) {
 }
 
 TEST(FxpTest, ExponentialPade) {
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = xt::linspace<float>(-22., 22., 4000);
 
@@ -79,7 +79,7 @@ TEST(FxpTest, ExponentialPade) {
 
 TEST(FxpTest, Log) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {{0.05, 0.5}, {5, 50}};
   // public log
@@ -110,7 +110,7 @@ TEST(FxpTest, Log) {
 
 TEST(FxpTest, Log2) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {{0.05, 0.5}, {5, 50}};
   // public log
@@ -141,7 +141,7 @@ TEST(FxpTest, Log2) {
 
 TEST(FxpTest, Log1p) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {{0.5, 2.0}, {0.9, 1.8}};
 
@@ -172,7 +172,7 @@ TEST(FxpTest, Log1p) {
 
 TEST(FxpTest, Exp2) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {{0.1, 0.2, 0.5, 0.7, 0.9},
                          {-0.1, -0.2, -0.5, -0.7, -0.9},
@@ -195,7 +195,7 @@ TEST(FxpTest, Exp2) {
 
 TEST(FxpTest, Tanh) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {{0.1, 0.2, 0.5, 0.7, 0.9, 2.1, 2.5, 4.0},
                          {-0.1, -0.2, -0.5, -0.7, -0.9, -2.1, -2.5, -4.0}};
@@ -220,7 +220,7 @@ TEST(FxpTest, Rsqrt) {
 
   // fxp_fraction_bits = 18(default value for FM64)
   {
-    HalContext ctx = test::makeRefHalContext();
+    SPUContext ctx = test::makeSPUContext();
 
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_rsqrt(&ctx, a);
@@ -238,7 +238,7 @@ TEST(FxpTest, Rsqrt) {
     config.set_protocol(ProtocolKind::REF2K);
     config.set_field(FieldType::FM64);
     config.set_fxp_fraction_bits(17);
-    HalContext ctx = test::makeRefHalContext(config);
+    SPUContext ctx = test::makeSPUContext(config, nullptr);
 
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_rsqrt(&ctx, a);
@@ -255,7 +255,7 @@ TEST(FxpTest, Rsqrt) {
     config.set_protocol(ProtocolKind::REF2K);
     config.set_field(FieldType::FM64);
     config.set_fxp_fraction_bits(16);
-    HalContext ctx = test::makeRefHalContext(config);
+    SPUContext ctx = test::makeSPUContext(config, nullptr);
 
     xt::random::seed(0);
     xt::xarray<float> x = xt::random::rand<float>({200, 1}, 256, 4096);
@@ -278,7 +278,7 @@ TEST(FxpTest, Sqrt) {
 
   // fxp_fraction_bits = 18(default value for FM64)
   {
-    HalContext ctx = test::makeRefHalContext();
+    SPUContext ctx = test::makeSPUContext();
 
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_sqrt(&ctx, a);
@@ -296,7 +296,7 @@ TEST(FxpTest, Sqrt) {
     config.set_protocol(ProtocolKind::REF2K);
     config.set_field(FieldType::FM64);
     config.set_fxp_fraction_bits(17);
-    HalContext ctx = test::makeRefHalContext(config);
+    SPUContext ctx = test::makeSPUContext(config, nullptr);
 
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_sqrt(&ctx, a);
diff --git a/libspu/kernel/hal/fxp_base_test.cc b/libspu/kernel/hal/fxp_base_test.cc
index 8df1db83..8ed1dad5 100644
--- a/libspu/kernel/hal/fxp_base_test.cc
+++ b/libspu/kernel/hal/fxp_base_test.cc
@@ -19,7 +19,7 @@
 
 #include "libspu/core/parallel_utils.h"
 #include "libspu/kernel/hal/constants.h"
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
 namespace {
@@ -67,7 +67,7 @@ INSTANTIATE_TEST_SUITE_P(
 // too slowly, disable for now. wait for linalg::matmul optimization.
 TEST_P(FxpMmulTest, Works) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   size_t m = std::get<0>(GetParam());
   size_t n = std::get<1>(GetParam());
@@ -102,7 +102,7 @@ TEST_P(FxpMmulTest, Works) {
 
 TEST(FxpTest, Reciprocal) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   // default fxp bits is 18 for FM64.
   xt::xarray<float> x = {
@@ -138,7 +138,7 @@ TEST(FxpTest, Reciprocal) {
 
 TEST(FxpTest, Div) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {{1.0, -200000.0, 7000000, -0.5, 314000, 1.5}};
   xt::xarray<float> y = {{1.0, 200000.0, 200000, 100, 3.14, 0.003}};
@@ -185,7 +185,7 @@ TEST(FxpTest, Div) {
 
 TEST(FxpTest, Abs) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {{0.5, -2.0}, {0.9, -1.8}};
 
@@ -217,7 +217,7 @@ TEST(FxpTest, Abs) {
 
 TEST(FxpTest, Floor) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {{0.5, -0.5}, {-20.0, 31.8}, {0, 5.0}, {-5.0, -31.8}};
 
@@ -249,7 +249,7 @@ TEST(FxpTest, Floor) {
 
 TEST(FxpTest, Ceil) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> x = {{0.5, -0.5}, {-20.0, 31.8}, {0, 5.0}, {-5.0, -31.8}};
 
diff --git a/libspu/kernel/hal/integer_test.cc b/libspu/kernel/hal/integer_test.cc
index 1da93f37..12fff27f 100644
--- a/libspu/kernel/hal/integer_test.cc
+++ b/libspu/kernel/hal/integer_test.cc
@@ -17,14 +17,14 @@
 #include "gtest/gtest.h"
 
 #include "libspu/kernel/hal/constants.h"
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
 
 // TODO: UT is too cumbersome.
 TEST(IntegralTest, Add) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   int ra = 3;
   int rb = 4;
@@ -50,7 +50,7 @@ TEST(IntegralTest, Add) {
 
 TEST(IntegralTest, Sub) {
   // GIVEN
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   int ra = 3;
   int rb = 4;
diff --git a/libspu/kernel/hal/polymorphic_test.cc b/libspu/kernel/hal/polymorphic_test.cc
index df04a927..7119d60e 100644
--- a/libspu/kernel/hal/polymorphic_test.cc
+++ b/libspu/kernel/hal/polymorphic_test.cc
@@ -23,7 +23,8 @@
 #include "xtensor/xmath.hpp"
 #include "xtensor/xvectorize.hpp"
 
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/hal/constants.h"
+#include "libspu/kernel/test_util.h"
 #include "libspu/mpc/utils/linalg.h"
 
 namespace spu::kernel::hal {
@@ -543,6 +544,7 @@ TYPED_TEST(MathUnaryTest, Log) {
   EXPECT_TRUE(xt::allclose(log_x, z, 0.1, 0.001)) << log_x << std::endl << z;
 }
 
+// TODO: can not pass MM1 & SEG3 test.
 TYPED_TEST(MathUnaryTest, Logistic) {
   using IN_DT = typename std::tuple_element<0, TypeParam>::type;
   using IN_VT = typename std::tuple_element<1, TypeParam>::type;
@@ -648,7 +650,7 @@ TEST_P(LogisticTest, Logistic) {
   config.set_protocol(ProtocolKind::REF2K);
   config.set_field(FieldType::FM64);
   config.set_sigmoid_mode(GetParam());
-  HalContext ctx = test::makeRefHalContext(config);
+  SPUContext ctx = test::makeSPUContext(config, nullptr);
 
   xt::xarray<float> x{{1.0, 2.0}, {0.5, 1.8}};
 
diff --git a/libspu/kernel/hal/ring_test.cc b/libspu/kernel/hal/ring_test.cc
index 9a9fbfe6..438fefbe 100644
--- a/libspu/kernel/hal/ring_test.cc
+++ b/libspu/kernel/hal/ring_test.cc
@@ -17,12 +17,12 @@
 #include "gtest/gtest.h"
 #include "xtensor/xarray.hpp"
 
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
 
 TEST(RingTest, _bit_parity) {
-  auto ctx = test::makeRefHalContext();
+  auto ctx = test::makeSPUContext();
   std::vector<std::pair<uint64_t, int>> para = {
       {0x0505, 0},
       {0x10505, 1},
@@ -42,7 +42,7 @@ TEST(RingTest, _bit_parity) {
 }
 
 TEST(RingTest, _popcount) {
-  auto ctx = test::makeRefHalContext();
+  auto ctx = test::makeSPUContext();
   std::vector<std::pair<uint64_t, int>> para = {
       {0x0505, 4},
       {0x10505, 5},
diff --git a/libspu/kernel/hal/shape_ops_test.cc b/libspu/kernel/hal/shape_ops_test.cc
index f26f7b8d..a58ec4ef 100644
--- a/libspu/kernel/hal/shape_ops_test.cc
+++ b/libspu/kernel/hal/shape_ops_test.cc
@@ -19,7 +19,7 @@
 #include "xtensor/xio.hpp"
 #include "xtensor/xshape.hpp"
 
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
 
@@ -48,7 +48,7 @@ TYPED_TEST(ShapeOpsUnaryTest, Transpose) {
   xt::xarray<IN_DT> x = test::xt_random<IN_DT>({2, 3, 4});
 
   // WHAT
-  auto transpose_wrapper = [](HalContext* ctx, const Value& x) {
+  auto transpose_wrapper = [](SPUContext* ctx, const Value& x) {
     return transpose(ctx, x);
   };
   auto z = test::evalUnaryOp<RES_DT>(IN_VT(), transpose_wrapper, x);
@@ -66,7 +66,7 @@ TYPED_TEST(ShapeOpsUnaryTest, TransposeWithPermutation) {
   // GIVEN
   xt::xarray<IN_DT> x = test::xt_random<IN_DT>({2, 2, 4});
 
-  auto transpose_wrapper = [](HalContext* ctx, const Value& x) {
+  auto transpose_wrapper = [](SPUContext* ctx, const Value& x) {
     return transpose(ctx, x, {1, 2, 0});
   };
 
@@ -87,7 +87,7 @@ TYPED_TEST(ShapeOpsUnaryTest, BroadcastTo) {
   // GIVEN
   xt::xarray<IN_DT> x = test::xt_random<IN_DT>({5, 1, 6});
 
-  auto broadcast_to_wrapper = [](HalContext* ctx, const Value& in) {
+  auto broadcast_to_wrapper = [](SPUContext* ctx, const Value& in) {
     return broadcast_to(ctx, in, {5, 4, 6});
   };
 
@@ -107,7 +107,7 @@ TYPED_TEST(ShapeOpsUnaryTest, BroadcastScalar) {
   // GIVEN
   xt::xarray<IN_DT> x = test::xt_random<IN_DT>({});
 
-  auto broadcast_to_wrapper = [](HalContext* ctx, const Value& in) {
+  auto broadcast_to_wrapper = [](SPUContext* ctx, const Value& in) {
     return broadcast_to(ctx, in, {1, 1});
   };
 
@@ -126,7 +126,7 @@ TYPED_TEST(ShapeOpsUnaryTest, BroadcastInDims) {
   // GIVEN
   xt::xarray<IN_DT> x = {1, 2, 3, 4};
 
-  auto broadcast_to_wrapper = [](HalContext* ctx, const Value& in) {
+  auto broadcast_to_wrapper = [](SPUContext* ctx, const Value& in) {
     return broadcast_to(ctx, in, {4, 2}, {0});
   };
 
@@ -144,7 +144,7 @@ TEST(SliceTest, Slice) {
   xt::xarray<int32_t> x = {{0, 1, 2}, {3, 4, 5}, {6, 7, 8}, {9, 10, 11}};
   using P_VT = public_v::type;
 
-  auto slice_wrapper = [](HalContext* ctx, const Value& in) {
+  auto slice_wrapper = [](SPUContext* ctx, const Value& in) {
     return slice(ctx, in, {2, 1}, {4, 3}, {});
   };
   auto z = test::evalUnaryOp<int64_t>(P_VT(), slice_wrapper, x);
@@ -161,7 +161,7 @@ TEST(SliceTest, UpdateSlice) {
   xt::xarray<int32_t> y = {{0, 1}, {3, 4}};
   using P_VT = public_v::type;
 
-  auto update_slice_wrapper = [](HalContext* ctx, const Value& in,
+  auto update_slice_wrapper = [](SPUContext* ctx, const Value& in,
                                  const Value& update) {
     return update_slice(ctx, in, update, {1, 1});
   };
@@ -176,7 +176,7 @@ TEST(SliceTest, SliceStride) {
   xt::xarray<int32_t> x = {0, 1, 2, 3};
   using P_VT = public_v::type;
 
-  auto slice_wrapper = [](HalContext* ctx, const Value& in) {
+  auto slice_wrapper = [](SPUContext* ctx, const Value& in) {
     return slice(ctx, in, {0}, {4}, {3});
   };
   auto z = test::evalUnaryOp<int64_t>(P_VT(), slice_wrapper, x);
@@ -193,7 +193,7 @@ TEST(ReshapeTest, Reshape) {
   xt::xarray<int32_t> x = {1, 2, 3, 4};
   using P_VT = public_v::type;
 
-  auto reshape_wrapper = [](HalContext* ctx, const Value& in) {
+  auto reshape_wrapper = [](SPUContext* ctx, const Value& in) {
     return reshape(ctx, in, {2, 2});
   };
   auto z = test::evalUnaryOp<int64_t>(P_VT(), reshape_wrapper, x);
@@ -210,7 +210,7 @@ TEST(ShapeOpsUnaryTest, Reverse) {
   };
   using P_VT = public_v::type;
 
-  auto reverse_wrapper = [](HalContext* ctx, const Value& in) {
+  auto reverse_wrapper = [](SPUContext* ctx, const Value& in) {
     return reverse(ctx, in, {0, 1, 2});
   };
 
@@ -231,7 +231,7 @@ TEST(ShapeOpsUnaryTest, BroadcastAfterReshape) {
   xt::xarray<int32_t> x = {{1, 2}, {3, 4}};
 
   using P_VT = public_v::type;
-  auto pad_wrapper = [](HalContext* ctx, const Value& in) {
+  auto pad_wrapper = [](SPUContext* ctx, const Value& in) {
     auto x1 = reshape(ctx, in, {1, 2, 2});
     return broadcast_to(ctx, x1, {2, 2, 2});
   };
@@ -253,7 +253,7 @@ TEST(ShapeOpsUnaryTest, Pad) {
   }}};
 
   using P_VT = public_v::type;
-  auto pad_wrapper = [](HalContext* ctx, const Value& in) {
+  auto pad_wrapper = [](SPUContext* ctx, const Value& in) {
     return pad(ctx, in, test::makeValue(ctx, 35, P_VT()), {1, 0, 0, 0},
                {0, 2, 0, 0}, {2, 1, 0, 0});
   };
@@ -280,7 +280,7 @@ TEST(ShapeOpsUnaryTest, InteriorPadding) {
                            {16, 17, 18, 19, 20}};
 
   using P_VT = public_v::type;
-  auto pad_wrapper = [](HalContext* ctx, const Value& in) {
+  auto pad_wrapper = [](SPUContext* ctx, const Value& in) {
     return pad(ctx, in, test::makeValue(ctx, 0, P_VT()), {0, 0}, {0, 0},
                {1, 1});
   };
@@ -307,7 +307,7 @@ TEST(ShapeOpsUnaryTest, NegativeEdgePad) {
                            {16, 17, 18, 19, 20}};
 
   using P_VT = public_v::type;
-  auto pad_wrapper = [](HalContext* ctx, const Value& in) {
+  auto pad_wrapper = [](SPUContext* ctx, const Value& in) {
     return pad(ctx, in, test::makeValue(ctx, 0, P_VT()), {-1, -1}, {-1, -1},
                {0, 0});
   };
@@ -329,7 +329,7 @@ TEST(ShapeOpsUnaryTest, NegativeEdgePadWithInteriorPad) {
                            {16, 17, 18, 19, 20}};
 
   using P_VT = public_v::type;
-  auto pad_wrapper = [](HalContext* ctx, const Value& in) {
+  auto pad_wrapper = [](SPUContext* ctx, const Value& in) {
     return pad(ctx, in, test::makeValue(ctx, 0, P_VT()), {-1, -1}, {-1, -1},
                {1, 1});
   };
@@ -354,7 +354,7 @@ TEST(ShapeOpsUnaryTest, HighNegativeEdgePadWithInteriorPad) {
                            {16, 17, 18, 19, 20}};
 
   using P_VT = public_v::type;
-  auto pad_wrapper = [](HalContext* ctx, const Value& in) {
+  auto pad_wrapper = [](SPUContext* ctx, const Value& in) {
     return pad(ctx, in, test::makeValue(ctx, 0, P_VT()), {-3, -3}, {-1, -1},
                {1, 1});
   };
@@ -398,7 +398,7 @@ TYPED_TEST(ConcatTest, Concatenate) {
   xt::xarray<DT> x = test::xt_random<DT>({3, 3});
   xt::xarray<DT> y = test::xt_random<DT>({3, 3});
 
-  auto concat_wrapper = [](HalContext* ctx, const Value& lhs,
+  auto concat_wrapper = [](SPUContext* ctx, const Value& lhs,
                            const Value& rhs) {
     return concatenate(ctx, {lhs, rhs}, 0);
   };
@@ -422,7 +422,7 @@ TYPED_TEST(ConcatTest, VConcatenate) {
   xt::xarray<DT> x = test::xt_random<DT>({3, 3});
   xt::xarray<DT> y = test::xt_random<DT>({3, 3});
 
-  auto concat_wrapper = [](HalContext* ctx, const Value& lhs,
+  auto concat_wrapper = [](SPUContext* ctx, const Value& lhs,
                            const Value& rhs) {
     return concatenate(ctx, {lhs, rhs}, 1);
   };
diff --git a/libspu/kernel/hal/test_util.cc b/libspu/kernel/hal/test_util.cc
deleted file mode 100644
index bfc00414..00000000
--- a/libspu/kernel/hal/test_util.cc
+++ /dev/null
@@ -1,55 +0,0 @@
-// Copyright 2021 Ant Group Co., Ltd.
-//
-// 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 "libspu/kernel/hal/test_util.h"
-
-#include "libspu/core/encoding.h"
-
-namespace spu::kernel::hal::test {
-
-HalContext makeRefHalContext(RuntimeConfig config) {
-  // Note: we are testing the encoding and approximation method, not the
-  // protocol itself, so ref2k is enough.
-  config.set_enable_action_trace(false);
-  HalContext ctx(config,  //
-                 nullptr  // link context.
-  );
-  return ctx;
-}
-
-HalContext makeRefHalContext() {
-  RuntimeConfig config;
-  config.set_protocol(ProtocolKind::REF2K);
-  config.set_field(FieldType::FM64);
-  config.set_sigmoid_mode(RuntimeConfig::SIGMOID_REAL);
-  return makeRefHalContext(config);
-}
-
-Value makeValue(HalContext* ctx, PtBufferView init, Visibility vtype,
-                DataType dtype, ShapeView shape) {
-  if (dtype == DT_INVALID) {
-    dtype = getEncodeType(init.pt_type);
-  }
-  auto res = constant(ctx, init, dtype, shape);
-  switch (vtype) {
-    case VIS_PUBLIC:
-      return res;
-    case VIS_SECRET:
-      return _p2s(ctx, res).setDtype(res.dtype());
-    default:
-      SPU_THROW("not supported vtype={}", vtype);
-  }
-}
-
-}  // namespace spu::kernel::hal::test
diff --git a/libspu/kernel/hal/type_cast_test.cc b/libspu/kernel/hal/type_cast_test.cc
index a824af46..247434e8 100644
--- a/libspu/kernel/hal/type_cast_test.cc
+++ b/libspu/kernel/hal/type_cast_test.cc
@@ -19,7 +19,8 @@
 #include "gtest/gtest.h"
 #include "xtensor/xio.hpp"
 
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/hal/constants.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
 namespace {
@@ -27,7 +28,7 @@ namespace {
 TEST(TypeCastTest, int2fxp) {
   const xt::xarray<int32_t> x = test::xt_random<int32_t>({5, 6});
   auto expected = xt::cast<int32_t>(x);
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   {
     Value a = constant(&ctx, x, DT_I32);
@@ -58,7 +59,7 @@ TEST(TypeCastTest, fxp2int) {
                                {-0.0, -1.0, -5.0, -10.0, -100.0, -1000.0},
                                {-0.1, -0.5, -0.7, -10.1, -10.5, -10.7}};
   auto expected = xt::cast<int32_t>(x);
-  HalContext ctx = test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
 
   {
     Value a = constant(&ctx, x, DT_FXP);
diff --git a/libspu/kernel/hlo/basic_binary_test.cc b/libspu/kernel/hlo/basic_binary_test.cc
index bba3c46c..6d00f18e 100644
--- a/libspu/kernel/hlo/basic_binary_test.cc
+++ b/libspu/kernel/hlo/basic_binary_test.cc
@@ -16,12 +16,13 @@
 
 #include "gtest/gtest.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/ndarray_ref.h"
-#include "libspu/kernel/context.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hlo/casting.h"
 #include "libspu/kernel/hlo/const.h"
-#include "libspu/kernel/hlo/test_utils.h"
-#include "libspu/kernel/value.h"
+#include "libspu/kernel/test_util.h"
+#include "libspu/mpc/factory.h"
 #include "libspu/mpc/utils/simulate.h"
 
 namespace spu::kernel::hlo {
@@ -29,20 +30,20 @@ namespace spu::kernel::hlo {
 class BinaryTest
     : public ::testing::TestWithParam<std::tuple<FieldType, ProtocolKind>> {};
 
-#define BINARY_EMPTY_TEST(NAME)                                                \
-  TEST_P(BinaryTest, Empty_##NAME) {                                           \
-    auto cfg =                                                                 \
-        test::makeRefConfig(std::get<0>(GetParam()), std::get<1>(GetParam())); \
-    mpc::utils::simulate(                                                      \
-        3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {             \
-          HalContext hctx(cfg, lctx);                                          \
-          auto empty_c0 = Seal(&hctx, Constant(&hctx, 1, {0}));                \
-          auto empty_c1 = Seal(&hctx, Constant(&hctx, 1, {0}));                \
-          auto s_empty = NAME(&hctx, empty_c0, empty_c1);                      \
-          EXPECT_EQ(s_empty.numel(), 0);                                       \
-          EXPECT_EQ(s_empty.shape().size(), 1);                                \
-          EXPECT_EQ(s_empty.shape()[0], 0);                                    \
-        });                                                                    \
+#define BINARY_EMPTY_TEST(NAME)                                      \
+  TEST_P(BinaryTest, Empty_##NAME) {                                 \
+    FieldType field = std::get<0>(GetParam());                       \
+    ProtocolKind prot = std::get<1>(GetParam());                     \
+    mpc::utils::simulate(                                            \
+        3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {   \
+          SPUContext sctx = test::makeSPUContext(prot, field, lctx); \
+          auto empty_c0 = Seal(&sctx, Constant(&sctx, 1, {0}));      \
+          auto empty_c1 = Seal(&sctx, Constant(&sctx, 1, {0}));      \
+          auto s_empty = NAME(&sctx, empty_c0, empty_c1);            \
+          EXPECT_EQ(s_empty.numel(), 0);                             \
+          EXPECT_EQ(s_empty.shape().size(), 1);                      \
+          EXPECT_EQ(s_empty.shape()[0], 0);                          \
+        });                                                          \
   }
 
 BINARY_EMPTY_TEST(Add)
@@ -64,17 +65,17 @@ BINARY_EMPTY_TEST(Div)
 BINARY_EMPTY_TEST(Remainder)
 
 TEST_P(BinaryTest, Empty_Dot) {
-  auto cfg =
-      test::makeRefConfig(std::get<0>(GetParam()), std::get<1>(GetParam()));
+  FieldType field = std::get<0>(GetParam());
+  ProtocolKind prot = std::get<1>(GetParam());
 
   mpc::utils::simulate(
       3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-        HalContext hctx(cfg, lctx);
-        auto empty_c0 = Seal(&hctx, Constant(&hctx, 1, {1, 0}));
-        auto empty_c1 = Seal(&hctx, Constant(&hctx, 1, {0, 1}));
+        SPUContext sctx = test::makeSPUContext(prot, field, lctx);
+        auto empty_c0 = Seal(&sctx, Constant(&sctx, 1, {1, 0}));
+        auto empty_c1 = Seal(&sctx, Constant(&sctx, 1, {0, 1}));
 
         // M = 1, N = 1, K = 0, result should be 1x1
-        auto ret = Dot(&hctx, empty_c0, empty_c1);
+        auto ret = Dot(&sctx, empty_c0, empty_c1);
 
         EXPECT_EQ(ret.numel(), 1);
         EXPECT_EQ(ret.shape().size(), 2);
@@ -82,7 +83,7 @@ TEST_P(BinaryTest, Empty_Dot) {
         EXPECT_EQ(ret.shape()[1], 1);
 
         // M = 0, N = 0, K = 1, result should be 0x0
-        ret = Dot(&hctx, empty_c1, empty_c0);
+        ret = Dot(&sctx, empty_c1, empty_c0);
 
         EXPECT_EQ(ret.numel(), 0);
         EXPECT_EQ(ret.shape().size(), 2);
@@ -92,17 +93,17 @@ TEST_P(BinaryTest, Empty_Dot) {
 }
 
 TEST_P(BinaryTest, MixedIntMul) {
-  auto cfg =
-      test::makeRefConfig(std::get<0>(GetParam()), std::get<1>(GetParam()));
+  FieldType field = std::get<0>(GetParam());
+  ProtocolKind prot = std::get<1>(GetParam());
 
   mpc::utils::simulate(
       3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-        HalContext hctx(cfg, lctx);
-        auto c0 = Constant(&hctx, static_cast<int32_t>(1), {1, 1});
-        auto c1 = Constant(&hctx, static_cast<int16_t>(1), {1, 1});
+        SPUContext sctx = test::makeSPUContext(prot, field, lctx);
+        auto c0 = Constant(&sctx, static_cast<int32_t>(1), {1, 1});
+        auto c1 = Constant(&sctx, static_cast<int16_t>(1), {1, 1});
 
         // M = 1, N = 1, K = 0, result should be 1x1
-        auto ret = Mul(&hctx, c0, c1);
+        auto ret = Mul(&sctx, c0, c1);
 
         EXPECT_EQ(ret.dtype(), DataType::DT_I32);
       });
diff --git a/libspu/kernel/hlo/basic_ternary_test.cc b/libspu/kernel/hlo/basic_ternary_test.cc
index c2d81c9c..64972edd 100644
--- a/libspu/kernel/hlo/basic_ternary_test.cc
+++ b/libspu/kernel/hlo/basic_ternary_test.cc
@@ -16,12 +16,12 @@
 
 #include "gtest/gtest.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/ndarray_ref.h"
-#include "libspu/kernel/context.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hlo/casting.h"
 #include "libspu/kernel/hlo/const.h"
-#include "libspu/kernel/hlo/test_utils.h"
-#include "libspu/kernel/value.h"
+#include "libspu/kernel/test_util.h"
 #include "libspu/mpc/utils/simulate.h"
 
 namespace spu::kernel::hlo {
@@ -30,17 +30,17 @@ class TernaryTest
     : public ::testing::TestWithParam<std::tuple<FieldType, ProtocolKind>> {};
 
 TEST_P(TernaryTest, SelectEmpty) {
-  auto cfg =
-      test::makeRefConfig(std::get<0>(GetParam()), std::get<1>(GetParam()));
+  FieldType field = std::get<0>(GetParam());
+  ProtocolKind prot = std::get<1>(GetParam());
 
   mpc::utils::simulate(
       3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-        HalContext hctx(cfg, lctx);
-        auto empty_p = Seal(&hctx, Constant(&hctx, true, {0}));
-        auto empty_true = Seal(&hctx, Constant(&hctx, 1, {0}));
-        auto empty_false = Seal(&hctx, Constant(&hctx, 2, {0}));
+        SPUContext sctx = test::makeSPUContext(prot, field, lctx);
+        auto empty_p = Seal(&sctx, Constant(&sctx, true, {0}));
+        auto empty_true = Seal(&sctx, Constant(&sctx, 1, {0}));
+        auto empty_false = Seal(&sctx, Constant(&sctx, 2, {0}));
 
-        auto ret = Select(&hctx, empty_p, empty_true, empty_false);
+        auto ret = Select(&sctx, empty_p, empty_true, empty_false);
 
         EXPECT_EQ(ret.numel(), 0);
         EXPECT_EQ(ret.shape().size(), 1);
@@ -49,17 +49,17 @@ TEST_P(TernaryTest, SelectEmpty) {
 }
 
 TEST_P(TernaryTest, ClampEmpty) {
-  auto cfg =
-      test::makeRefConfig(std::get<0>(GetParam()), std::get<1>(GetParam()));
+  FieldType field = std::get<0>(GetParam());
+  ProtocolKind prot = std::get<1>(GetParam());
 
   mpc::utils::simulate(
       3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-        HalContext hctx(cfg, lctx);
-        auto empty_in = Seal(&hctx, Constant(&hctx, 0, {0}));
-        auto empty_min = Seal(&hctx, Constant(&hctx, 1, {0}));
-        auto empty_max = Seal(&hctx, Constant(&hctx, 2, {0}));
+        SPUContext sctx = test::makeSPUContext(prot, field, lctx);
+        auto empty_in = Seal(&sctx, Constant(&sctx, 0, {0}));
+        auto empty_min = Seal(&sctx, Constant(&sctx, 1, {0}));
+        auto empty_max = Seal(&sctx, Constant(&sctx, 2, {0}));
 
-        auto ret = Clamp(&hctx, empty_in, empty_min, empty_max);
+        auto ret = Clamp(&sctx, empty_in, empty_min, empty_max);
 
         EXPECT_EQ(ret.numel(), 0);
         EXPECT_EQ(ret.shape().size(), 1);
diff --git a/libspu/kernel/hlo/basic_unary_test.cc b/libspu/kernel/hlo/basic_unary_test.cc
index a2eade0c..677cb204 100644
--- a/libspu/kernel/hlo/basic_unary_test.cc
+++ b/libspu/kernel/hlo/basic_unary_test.cc
@@ -16,12 +16,12 @@
 
 #include "gtest/gtest.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/ndarray_ref.h"
-#include "libspu/kernel/context.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hlo/casting.h"
 #include "libspu/kernel/hlo/const.h"
-#include "libspu/kernel/hlo/test_utils.h"
-#include "libspu/kernel/value.h"
+#include "libspu/kernel/test_util.h"
 #include "libspu/mpc/utils/simulate.h"
 
 namespace spu::kernel::hlo {
@@ -29,19 +29,19 @@ namespace spu::kernel::hlo {
 class UnaryTest
     : public ::testing::TestWithParam<std::tuple<FieldType, ProtocolKind>> {};
 
-#define UNARY_EMPTY_TEST(NAME)                                                 \
-  TEST_P(UnaryTest, Empty_##NAME) {                                            \
-    auto cfg =                                                                 \
-        test::makeRefConfig(std::get<0>(GetParam()), std::get<1>(GetParam())); \
-    mpc::utils::simulate(                                                      \
-        3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {             \
-          HalContext hctx(cfg, lctx);                                          \
-          auto empty_c = Constant(&hctx, 1.0F, {0});                           \
-          auto s_empty = NAME(&hctx, empty_c);                                 \
-          EXPECT_EQ(s_empty.numel(), 0);                                       \
-          EXPECT_EQ(s_empty.shape().size(), 1);                                \
-          EXPECT_EQ(s_empty.shape()[0], 0);                                    \
-        });                                                                    \
+#define UNARY_EMPTY_TEST(NAME)                                       \
+  TEST_P(UnaryTest, Empty_##NAME) {                                  \
+    FieldType field = std::get<0>(GetParam());                       \
+    ProtocolKind prot = std::get<1>(GetParam());                     \
+    mpc::utils::simulate(                                            \
+        3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {   \
+          SPUContext sctx = test::makeSPUContext(prot, field, lctx); \
+          auto empty_c = Constant(&sctx, 1.0F, {0});                 \
+          auto s_empty = NAME(&sctx, empty_c);                       \
+          EXPECT_EQ(s_empty.numel(), 0);                             \
+          EXPECT_EQ(s_empty.shape().size(), 1);                      \
+          EXPECT_EQ(s_empty.shape()[0], 0);                          \
+        });                                                          \
   }
 
 UNARY_EMPTY_TEST(Reciprocal)
diff --git a/libspu/kernel/hlo/casting_test.cc b/libspu/kernel/hlo/casting_test.cc
index 8cf42db2..f103a155 100644
--- a/libspu/kernel/hlo/casting_test.cc
+++ b/libspu/kernel/hlo/casting_test.cc
@@ -16,23 +16,23 @@
 
 #include "gtest/gtest.h"
 
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/core/context.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hlo/const.h"
-#include "libspu/kernel/value.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hlo {
 
 TEST(ConstTest, Empty) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
 
-  auto empty_c = Constant(&hctx, true, {0});
+  auto empty_c = Constant(&sctx, true, {0});
 
   // Seal
-  auto s_empty = Seal(&hctx, empty_c);
+  auto s_empty = Seal(&sctx, empty_c);
 
   // Reveal
-  auto p_empty = Reveal(&hctx, s_empty);
+  auto p_empty = Reveal(&sctx, s_empty);
 
   EXPECT_EQ(p_empty.numel(), 0);
   EXPECT_EQ(p_empty.shape().size(), 1);
diff --git a/libspu/kernel/hlo/const_test.cc b/libspu/kernel/hlo/const_test.cc
index ab90fa14..17220169 100644
--- a/libspu/kernel/hlo/const_test.cc
+++ b/libspu/kernel/hlo/const_test.cc
@@ -17,17 +17,17 @@
 #include "gtest/gtest.h"
 #include "spdlog/fmt/bin_to_hex.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/ndarray_ref.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/hal/test_util.h"
-#include "libspu/kernel/value.h"
+#include "libspu/core/value.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hlo {
 
 TEST(ConstTest, Empty) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
 
-  auto empty_c = Constant(&hctx, true, {0});
+  auto empty_c = Constant(&sctx, true, {0});
 
   EXPECT_EQ(empty_c.numel(), 0);
   EXPECT_EQ(empty_c.shape().size(), 1);
@@ -35,14 +35,13 @@ TEST(ConstTest, Empty) {
 }
 
 TEST(ConstTest, Epsilon) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
 
-  auto eps = Epsilon(&hctx);
+  auto eps = Epsilon(&sctx);
 
-  auto v = hal::dump_public_as<float>(&hctx, eps);
+  auto v = hal::dump_public_as<float>(&sctx, eps);
 
-  EXPECT_FLOAT_EQ(v[0],
-                  1 / (std::pow(2, hctx.rt_config().fxp_fraction_bits())));
+  EXPECT_FLOAT_EQ(v[0], 1 / (std::pow(2, sctx.config().fxp_fraction_bits())));
 }
 
 }  // namespace spu::kernel::hlo
diff --git a/libspu/kernel/hlo/geometrical_test.cc b/libspu/kernel/hlo/geometrical_test.cc
index 8c56ddda..be1f5f3c 100644
--- a/libspu/kernel/hlo/geometrical_test.cc
+++ b/libspu/kernel/hlo/geometrical_test.cc
@@ -16,12 +16,12 @@
 
 #include "gtest/gtest.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/ndarray_ref.h"
-#include "libspu/kernel/context.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hlo/basic_binary.h"
 #include "libspu/kernel/hlo/const.h"
-#include "libspu/kernel/hlo/test_utils.h"
-#include "libspu/kernel/value.h"
+#include "libspu/kernel/test_util.h"
 #include "libspu/mpc/utils/simulate.h"
 
 namespace spu::kernel::hlo {
@@ -30,40 +30,40 @@ class GeoTest
     : public ::testing::TestWithParam<std::tuple<FieldType, ProtocolKind>> {};
 
 TEST_P(GeoTest, EmptySlice) {
-  auto cfg =
-      test::makeRefConfig(std::get<0>(GetParam()), std::get<1>(GetParam()));
+  FieldType field = std::get<0>(GetParam());
+  ProtocolKind prot = std::get<1>(GetParam());
 
-  mpc::utils::simulate(3,
-                       [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-                         HalContext hctx(cfg, lctx);
-                         auto in = Iota(&hctx, DT_I64, 10);
+  mpc::utils::simulate(
+      3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
+        SPUContext sctx = test::makeSPUContext(prot, field, lctx);
+        auto in = Iota(&sctx, DT_I64, 10);
 
-                         auto ret = Slice(&hctx, in, {1}, {1}, {1});
+        auto ret = Slice(&sctx, in, {1}, {1}, {1});
 
-                         EXPECT_EQ(ret.numel(), 0);
-                         EXPECT_EQ(ret.shape().size(), 1);
-                         EXPECT_EQ(ret.shape()[0], 0);
+        EXPECT_EQ(ret.numel(), 0);
+        EXPECT_EQ(ret.shape().size(), 1);
+        EXPECT_EQ(ret.shape()[0], 0);
 
-                         ret = Slice(&hctx, in, {2}, {1}, {1});
-                         EXPECT_EQ(ret.numel(), 0);
-                         EXPECT_EQ(ret.shape().size(), 1);
-                         EXPECT_EQ(ret.shape()[0], 0);
-                       });
+        ret = Slice(&sctx, in, {2}, {1}, {1});
+        EXPECT_EQ(ret.numel(), 0);
+        EXPECT_EQ(ret.shape().size(), 1);
+        EXPECT_EQ(ret.shape()[0], 0);
+      });
 }
 
 TEST_P(GeoTest, MixedStoragePad) {
-  auto cfg =
-      test::makeRefConfig(std::get<0>(GetParam()), std::get<1>(GetParam()));
+  FieldType field = std::get<0>(GetParam());
+  ProtocolKind prot = std::get<1>(GetParam());
 
   mpc::utils::simulate(
       3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-        HalContext hctx(cfg, lctx);
-        auto in = Iota(&hctx, DT_I64, 10);
-        in = Seal(&hctx, in);
-        auto pv = Constant(&hctx, static_cast<int64_t>(0), {1});
-        pv = Xor(&hctx, pv, pv);  // Force a bshr
+        SPUContext sctx = test::makeSPUContext(prot, field, lctx);
+        auto in = Iota(&sctx, DT_I64, 10);
+        in = Seal(&sctx, in);
+        auto pv = Constant(&sctx, static_cast<int64_t>(0), {1});
+        pv = Xor(&sctx, pv, pv);  // Force a bshr
 
-        auto ret = Pad(&hctx, in, pv, {1}, {0}, {0});
+        auto ret = Pad(&sctx, in, pv, {1}, {0}, {0});
 
         EXPECT_EQ(ret.numel(), 11);
         EXPECT_EQ(ret.vtype(), VIS_SECRET);
@@ -72,20 +72,20 @@ TEST_P(GeoTest, MixedStoragePad) {
 }
 
 TEST_P(GeoTest, Pad) {
-  auto cfg =
-      test::makeRefConfig(std::get<0>(GetParam()), std::get<1>(GetParam()));
+  FieldType field = std::get<0>(GetParam());
+  ProtocolKind prot = std::get<1>(GetParam());
 
   mpc::utils::simulate(
       3, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-        HalContext hctx(cfg, lctx);
+        SPUContext sctx = test::makeSPUContext(prot, field, lctx);
         auto in = Seal(
-            &hctx, Cast(&hctx, Constant(&hctx, std::vector<uint8_t>(2, 0), {2}),
+            &sctx, Cast(&sctx, Constant(&sctx, std::vector<uint8_t>(2, 0), {2}),
                         VIS_PUBLIC, DT_I1));
         auto pv = Seal(
-            &hctx, Cast(&hctx, Constant(&hctx, static_cast<uint8_t>(0), {1}),
+            &sctx, Cast(&sctx, Constant(&sctx, static_cast<uint8_t>(0), {1}),
                         VIS_PUBLIC, DT_I1));
 
-        auto ret = Pad(&hctx, in, pv, {0}, {1}, {0});
+        auto ret = Pad(&sctx, in, pv, {0}, {1}, {0});
 
         EXPECT_EQ(ret.numel(), 3);
         EXPECT_EQ(ret.vtype(), VIS_SECRET);
diff --git a/libspu/kernel/hlo/indexing_test.cc b/libspu/kernel/hlo/indexing_test.cc
index 32618736..6cf54232 100644
--- a/libspu/kernel/hlo/indexing_test.cc
+++ b/libspu/kernel/hlo/indexing_test.cc
@@ -16,12 +16,12 @@
 
 #include "gtest/gtest.h"
 
+#include "libspu/core/context.h"
 #include "libspu/core/ndarray_ref.h"
 #include "libspu/core/type.h"
-#include "libspu/kernel/context.h"
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/core/value.h"
 #include "libspu/kernel/hlo/const.h"
-#include "libspu/kernel/value.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hlo {
 
@@ -66,81 +66,81 @@ TEST(IndexingTest, Take2) {
 }
 
 TEST(IndexingTest, DynamicUpdateSliceScalarWithPublicIndices) {
-  HalContext hctx = hal::test::makeRefHalContext();
-  auto input = Constant(&hctx, static_cast<int64_t>(1), {5});
-  auto update = Constant(&hctx, static_cast<int64_t>(2), {1});
+  SPUContext sctx = test::makeSPUContext();
+  auto input = Constant(&sctx, static_cast<int64_t>(1), {5});
+  auto update = Constant(&sctx, static_cast<int64_t>(2), {1});
   std::vector<spu::Value> start_indices{
-      Constant(&hctx, static_cast<int64_t>(1), {})};
+      Constant(&sctx, static_cast<int64_t>(1), {})};
 
-  auto output = DynamicUpdateSlice(&hctx, input, update, start_indices);
+  auto output = DynamicUpdateSlice(&sctx, input, update, start_indices);
 
   EXPECT_EQ(output.numel(), input.numel());
 
-  auto p_ret = hal::dump_public_as<int64_t>(&hctx, output);
+  auto p_ret = hal::dump_public_as<int64_t>(&sctx, output);
   xt::xarray<int64_t> expected{1, 2, 1, 1, 1};
   EXPECT_EQ(p_ret, expected);
 }
 
 TEST(IndexingTest, DynamicUpdateSliceScalarWithSecretIndices) {
-  HalContext hctx = hal::test::makeRefHalContext();
-  auto input = Constant(&hctx, static_cast<int64_t>(1), {5});
-  auto update = Constant(&hctx, static_cast<int64_t>(2), {1});
+  SPUContext sctx = test::makeSPUContext();
+  auto input = Constant(&sctx, static_cast<int64_t>(1), {5});
+  auto update = Constant(&sctx, static_cast<int64_t>(2), {1});
   std::vector<spu::Value> start_indices{
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(1), {}))};
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(1), {}))};
 
-  auto output = DynamicUpdateSlice(&hctx, input, update, start_indices);
+  auto output = DynamicUpdateSlice(&sctx, input, update, start_indices);
 
   EXPECT_EQ(output.numel(), input.numel());
 
-  auto p_ret = hal::dump_public_as<int64_t>(&hctx, Reveal(&hctx, output));
+  auto p_ret = hal::dump_public_as<int64_t>(&sctx, Reveal(&sctx, output));
   xt::xarray<int64_t> expected{1, 2, 1, 1, 1};
   EXPECT_EQ(p_ret, expected);
 }
 
 TEST(IndexingTest, DynamicUpdateSliceMultiValuesWithPublicIndices) {
-  HalContext hctx = hal::test::makeRefHalContext();
-  auto input = Constant(&hctx, static_cast<int64_t>(1), {5});
-  auto update = Constant(&hctx, static_cast<int64_t>(2), {2});
+  SPUContext sctx = test::makeSPUContext();
+  auto input = Constant(&sctx, static_cast<int64_t>(1), {5});
+  auto update = Constant(&sctx, static_cast<int64_t>(2), {2});
   std::vector<spu::Value> start_indices{
-      Constant(&hctx, static_cast<int64_t>(3), {})};
+      Constant(&sctx, static_cast<int64_t>(3), {})};
 
-  auto output = DynamicUpdateSlice(&hctx, input, update, start_indices);
+  auto output = DynamicUpdateSlice(&sctx, input, update, start_indices);
 
   EXPECT_EQ(output.numel(), input.numel());
 
-  auto p_ret = hal::dump_public_as<int64_t>(&hctx, output);
+  auto p_ret = hal::dump_public_as<int64_t>(&sctx, output);
   xt::xarray<int64_t> expected{1, 1, 1, 2, 2};
   EXPECT_EQ(p_ret, expected);
 }
 
 TEST(IndexingTest, DynamicUpdateSliceMultiValuesWithSecretIndices) {
-  HalContext hctx = hal::test::makeRefHalContext();
-  auto input = Constant(&hctx, static_cast<int64_t>(1), {5});
-  auto update = Constant(&hctx, static_cast<int64_t>(2), {2});
+  SPUContext sctx = test::makeSPUContext();
+  auto input = Constant(&sctx, static_cast<int64_t>(1), {5});
+  auto update = Constant(&sctx, static_cast<int64_t>(2), {2});
   std::vector<spu::Value> start_indices{
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(3), {}))};
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(3), {}))};
 
-  auto output = DynamicUpdateSlice(&hctx, input, update, start_indices);
+  auto output = DynamicUpdateSlice(&sctx, input, update, start_indices);
 
   EXPECT_EQ(output.numel(), input.numel());
 
-  auto p_ret = hal::dump_public_as<int64_t>(&hctx, Reveal(&hctx, output));
+  auto p_ret = hal::dump_public_as<int64_t>(&sctx, Reveal(&sctx, output));
   xt::xarray<int64_t> expected{1, 1, 1, 2, 2};
   EXPECT_EQ(p_ret, expected);
 }
 
 TEST(DynamicSliceTest, DynamicSliceWithPublicIndices) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
   xt::xarray<float> x = {{0.05, 0.24, 0.5}, {2, 5, 50}};
-  auto input = hal::test::makeValue(&hctx, x, VIS_SECRET);
+  auto input = test::makeValue(&sctx, x, VIS_SECRET);
 
   auto start_indices =
-      std::vector<spu::Value>{Constant(&hctx, static_cast<int64_t>(1), {}),
-                              Constant(&hctx, static_cast<int64_t>(1), {})};
+      std::vector<spu::Value>{Constant(&sctx, static_cast<int64_t>(1), {}),
+                              Constant(&sctx, static_cast<int64_t>(1), {})};
 
-  auto output = DynamicSlice(&hctx, input, {2, 2}, start_indices);
+  auto output = DynamicSlice(&sctx, input, {2, 2}, start_indices);
 
-  auto p_ret = hal::dump_public_as<float>(&hctx, Reveal(&hctx, output));
+  auto p_ret = hal::dump_public_as<float>(&sctx, Reveal(&sctx, output));
   xt::xarray<float> expected{{0.24, 0.5}, {5, 50}};
   EXPECT_TRUE(xt::allclose(p_ret, expected, 0.01, 0.001))
       << p_ret << std::endl
@@ -148,17 +148,17 @@ TEST(DynamicSliceTest, DynamicSliceWithPublicIndices) {
 }
 
 TEST(DynamicSliceTest, DynamicSliceWithSecretIndices) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
   xt::xarray<float> x = {{0.05, 0.24, 0.5}, {2, 5, 50}, {7, 9, 10.1}};
-  auto input = hal::test::makeValue(&hctx, x, VIS_SECRET);
+  auto input = test::makeValue(&sctx, x, VIS_SECRET);
 
   auto start_indices = std::vector<spu::Value>{
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(0), {})),
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(1), {}))};
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(0), {})),
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(1), {}))};
 
-  auto output = DynamicSlice(&hctx, input, {2, 2}, start_indices);
+  auto output = DynamicSlice(&sctx, input, {2, 2}, start_indices);
 
-  auto p_ret = hal::dump_public_as<float>(&hctx, Reveal(&hctx, output));
+  auto p_ret = hal::dump_public_as<float>(&sctx, Reveal(&sctx, output));
   xt::xarray<float> expected{{0.24, 0.5}, {5, 50}};
   EXPECT_TRUE(xt::allclose(p_ret, expected, 0.01, 0.001))
       << p_ret << std::endl
@@ -166,17 +166,17 @@ TEST(DynamicSliceTest, DynamicSliceWithSecretIndices) {
 }
 
 TEST(DynamicSliceTest, DynamicSliceWithSecretIndicesFull) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
   xt::xarray<float> x = {{0.05, 0.24, 0.5}, {2, 5, 50}, {7, 9, 10.1}};
-  auto input = hal::test::makeValue(&hctx, x, VIS_SECRET);
+  auto input = test::makeValue(&sctx, x, VIS_SECRET);
 
   auto start_indices = std::vector<spu::Value>{
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(0), {})),
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(0), {}))};
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(0), {})),
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(0), {}))};
 
-  auto output = DynamicSlice(&hctx, input, {2, 3}, start_indices);
+  auto output = DynamicSlice(&sctx, input, {2, 3}, start_indices);
 
-  auto p_ret = hal::dump_public_as<float>(&hctx, Reveal(&hctx, output));
+  auto p_ret = hal::dump_public_as<float>(&sctx, Reveal(&sctx, output));
   xt::xarray<float> expected{{0.05, 0.24, 0.5}, {2, 5, 50}};
   EXPECT_TRUE(xt::allclose(p_ret, expected, 0.01, 0.001))
       << p_ret << std::endl
@@ -184,17 +184,17 @@ TEST(DynamicSliceTest, DynamicSliceWithSecretIndicesFull) {
 }
 
 TEST(DynamicSliceTest, DynamicSliceWithPublicIndicesOffRangeLow) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
   xt::xarray<float> x = {{0.05, 0.24, 0.5}, {2, 5, 50}};
-  auto input = hal::test::makeValue(&hctx, x, VIS_SECRET);
+  auto input = test::makeValue(&sctx, x, VIS_SECRET);
 
   auto start_indices =
-      std::vector<spu::Value>{Constant(&hctx, static_cast<int64_t>(-1), {}),
-                              Constant(&hctx, static_cast<int64_t>(-1), {})};
+      std::vector<spu::Value>{Constant(&sctx, static_cast<int64_t>(-1), {}),
+                              Constant(&sctx, static_cast<int64_t>(-1), {})};
 
-  auto output = DynamicSlice(&hctx, input, {2, 2}, start_indices);
+  auto output = DynamicSlice(&sctx, input, {2, 2}, start_indices);
 
-  auto p_ret = hal::dump_public_as<float>(&hctx, Reveal(&hctx, output));
+  auto p_ret = hal::dump_public_as<float>(&sctx, Reveal(&sctx, output));
   xt::xarray<float> expected{{0.05, 0.24}, {2, 5}};
   EXPECT_TRUE(xt::allclose(p_ret, expected, 0.01, 0.001))
       << p_ret << std::endl
@@ -202,17 +202,17 @@ TEST(DynamicSliceTest, DynamicSliceWithPublicIndicesOffRangeLow) {
 }
 
 TEST(DynamicSliceTest, DynamicSliceWithPublicIndicesOffRangeHigh) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
   xt::xarray<float> x = {{0.05, 0.24, 0.5}, {2, 5, 50}};
-  auto input = hal::test::makeValue(&hctx, x, VIS_SECRET);
+  auto input = test::makeValue(&sctx, x, VIS_SECRET);
 
   auto start_indices =
-      std::vector<spu::Value>{Constant(&hctx, static_cast<int64_t>(10), {}),
-                              Constant(&hctx, static_cast<int64_t>(10), {})};
+      std::vector<spu::Value>{Constant(&sctx, static_cast<int64_t>(10), {}),
+                              Constant(&sctx, static_cast<int64_t>(10), {})};
 
-  auto output = DynamicSlice(&hctx, input, {2, 2}, start_indices);
+  auto output = DynamicSlice(&sctx, input, {2, 2}, start_indices);
 
-  auto p_ret = hal::dump_public_as<float>(&hctx, Reveal(&hctx, output));
+  auto p_ret = hal::dump_public_as<float>(&sctx, Reveal(&sctx, output));
   xt::xarray<float> expected{{0.24, 0.5}, {5, 50}};
   EXPECT_TRUE(xt::allclose(p_ret, expected, 0.01, 0.001))
       << p_ret << std::endl
@@ -220,17 +220,17 @@ TEST(DynamicSliceTest, DynamicSliceWithPublicIndicesOffRangeHigh) {
 }
 
 TEST(DynamicSliceTest, DynamicSliceWithSecretIndicesOffRangeLow) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
   xt::xarray<float> x = {{0.05, 0.24, 0.5}, {2, 5, 50}};
-  auto input = hal::test::makeValue(&hctx, x, VIS_SECRET);
+  auto input = test::makeValue(&sctx, x, VIS_SECRET);
 
   auto start_indices = std::vector<spu::Value>{
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(-1), {})),
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(-1), {}))};
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(-1), {})),
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(-1), {}))};
 
-  auto output = DynamicSlice(&hctx, input, {2, 2}, start_indices);
+  auto output = DynamicSlice(&sctx, input, {2, 2}, start_indices);
 
-  auto p_ret = hal::dump_public_as<float>(&hctx, Reveal(&hctx, output));
+  auto p_ret = hal::dump_public_as<float>(&sctx, Reveal(&sctx, output));
   xt::xarray<float> expected{{0.05, 0.24}, {2, 5}};
   EXPECT_TRUE(xt::allclose(p_ret, expected, 0.01, 0.001))
       << p_ret << std::endl
@@ -238,17 +238,17 @@ TEST(DynamicSliceTest, DynamicSliceWithSecretIndicesOffRangeLow) {
 }
 
 TEST(DynamicSliceTest, DynamicSliceWithSecretIndicesOffRangeHigh) {
-  HalContext hctx = hal::test::makeRefHalContext();
+  SPUContext sctx = test::makeSPUContext();
   xt::xarray<float> x = {{0.05, 0.24, 0.5}, {2, 5, 50}};
-  auto input = hal::test::makeValue(&hctx, x, VIS_SECRET);
+  auto input = test::makeValue(&sctx, x, VIS_SECRET);
 
   auto start_indices = std::vector<spu::Value>{
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(10), {})),
-      Seal(&hctx, Constant(&hctx, static_cast<int64_t>(10), {}))};
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(10), {})),
+      Seal(&sctx, Constant(&sctx, static_cast<int64_t>(10), {}))};
 
-  auto output = DynamicSlice(&hctx, input, {2, 2}, start_indices);
+  auto output = DynamicSlice(&sctx, input, {2, 2}, start_indices);
 
-  auto p_ret = hal::dump_public_as<float>(&hctx, Reveal(&hctx, output));
+  auto p_ret = hal::dump_public_as<float>(&sctx, Reveal(&sctx, output));
   xt::xarray<float> expected{{0.24, 0.5}, {5, 50}};
   EXPECT_TRUE(xt::allclose(p_ret, expected, 0.01, 0.001))
       << p_ret << std::endl
diff --git a/libspu/kernel/hlo/shuffle_test.cc b/libspu/kernel/hlo/shuffle_test.cc
index 08b4afc6..0550ab58 100644
--- a/libspu/kernel/hlo/shuffle_test.cc
+++ b/libspu/kernel/hlo/shuffle_test.cc
@@ -18,14 +18,14 @@
 #include "xtensor/xio.hpp"
 #include "xtensor/xsort.hpp"
 
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hlo {
 
 TEST(SortTest, Array) {
-  HalContext ctx = hal::test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
   xt::xarray<float> x = xt::random::rand<float>({10});
-  std::vector<Value> x_v = {hal::test::makeValue(&ctx, x, VIS_SECRET)};
+  std::vector<Value> x_v = {test::makeValue(&ctx, x, VIS_SECRET)};
   spu::Value ret1 = Shuffle(&ctx, x_v, 0)[0];
   auto ret1_hat =
       hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, ret1).asFxp());
@@ -45,9 +45,9 @@ TEST(SortTest, Array) {
 }
 
 TEST(SortTest, 2D) {
-  HalContext ctx = hal::test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
   xt::xarray<float> x = xt::random::rand<float>({10, 15});
-  std::vector<Value> x_v = {hal::test::makeValue(&ctx, x, VIS_SECRET)};
+  std::vector<Value> x_v = {test::makeValue(&ctx, x, VIS_SECRET)};
 
   spu::Value ret1 = Shuffle(&ctx, x_v, 1)[0];
   auto ret1_hat =
diff --git a/libspu/kernel/hlo/sort_test.cc b/libspu/kernel/hlo/sort_test.cc
index 43ec7c3b..874a92db 100644
--- a/libspu/kernel/hlo/sort_test.cc
+++ b/libspu/kernel/hlo/sort_test.cc
@@ -21,16 +21,16 @@
 
 #include "libspu/kernel/hal/constants.h"
 #include "libspu/kernel/hal/polymorphic.h"
-#include "libspu/kernel/hal/test_util.h"
+#include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hlo {
 
 TEST(SortTest, Simple) {
-  HalContext ctx = hal::test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
   xt::xarray<float> x = {{0.05, 0.5, 0.24}, {5, 50, 2}};
   xt::xarray<float> sorted_x = {{0.05, 0.24, 0.5}, {2, 5, 50}};
 
-  Value x_v = hal::test::makeValue(&ctx, x, VIS_SECRET);
+  Value x_v = test::makeValue(&ctx, x, VIS_SECRET);
 
   std::vector<spu::Value> rets = Sort(
       &ctx, {x_v}, 1, false,
@@ -50,11 +50,11 @@ TEST(SortTest, Simple) {
 }
 
 TEST(SortTest, SimpleWithNoPadding) {
-  HalContext ctx = hal::test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
   xt::xarray<float> x = {{0.05, 0.5, 0.24, 0.5}, {2, 5, 50, 2}};
   xt::xarray<float> sorted_x = {{0.05, 0.24, 0.5, 0.5}, {2, 2, 5, 50}};
 
-  Value x_v = hal::test::makeValue(&ctx, x, VIS_SECRET);
+  Value x_v = test::makeValue(&ctx, x, VIS_SECRET);
 
   std::vector<spu::Value> rets = Sort(
       &ctx, {x_v}, 1, false,
@@ -74,14 +74,14 @@ TEST(SortTest, SimpleWithNoPadding) {
 }
 
 TEST(SortTest, MultiInputs) {
-  HalContext ctx = hal::test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
   xt::xarray<float> x1 = {{0.5, 0.05, 0.5, 0.24, 0.5, 0.5, 0.5}};
   xt::xarray<float> x2 = {{5, 1, 2, 1, 2, 3, 4}};
   xt::xarray<float> sorted_x1 = {{0.05, 0.24, 0.5, 0.5, 0.5, 0.5, 0.5}};
   xt::xarray<float> sorted_x2 = {{1, 1, 2, 2, 3, 4, 5}};
 
-  Value x1_v = hal::test::makeValue(&ctx, x1, VIS_SECRET);
-  Value x2_v = hal::test::makeValue(&ctx, x2, VIS_SECRET);
+  Value x1_v = test::makeValue(&ctx, x1, VIS_SECRET);
+  Value x2_v = test::makeValue(&ctx, x2, VIS_SECRET);
 
   std::vector<spu::Value> rets = Sort(
       &ctx, {x1_v, x2_v}, 1, false,
@@ -109,15 +109,15 @@ TEST(SortTest, MultiInputs) {
 }
 
 TEST(SortTest, MultiOperands) {
-  HalContext ctx = hal::test::makeRefHalContext();
+  SPUContext ctx = test::makeSPUContext();
   xt::xarray<float> k1 = {6, 6, 3, 4, 4, 5, 4};
   xt::xarray<float> k2 = {0.5, 0.1, 3.1, 6.5, 4.1, 6.7, 2.5};
 
   xt::xarray<float> sorted_k1 = {3, 4, 4, 4, 5, 6, 6};
   xt::xarray<float> sorted_k2 = {3.1, 2.5, 4.1, 6.5, 6.7, 0.1, 0.5};
 
-  Value k1_v = hal::test::makeValue(&ctx, k1, VIS_SECRET);
-  Value k2_v = hal::test::makeValue(&ctx, k2, VIS_SECRET);
+  Value k1_v = test::makeValue(&ctx, k1, VIS_SECRET);
+  Value k2_v = test::makeValue(&ctx, k2, VIS_SECRET);
 
   std::vector<spu::Value> rets = Sort(
       &ctx, {k1_v, k2_v}, 0, false,
diff --git a/libspu/kernel/hlo/test_utils.h b/libspu/kernel/hlo/test_utils.h
deleted file mode 100644
index 2c25d796..00000000
--- a/libspu/kernel/hlo/test_utils.h
+++ /dev/null
@@ -1,28 +0,0 @@
-// Copyright 2023 Ant Group Co., Ltd.
-//
-// 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.
-
-#pragma once
-
-#include "libspu/spu.pb.h"
-
-namespace spu::kernel::hlo::test {
-
-inline auto makeRefConfig(FieldType ft, ProtocolKind pk) {
-  RuntimeConfig cfg;
-  cfg.set_protocol(pk);
-  cfg.set_field(ft);
-  return cfg;
-}
-
-}  // namespace spu::kernel::hlo::test
diff --git a/libspu/kernel/test_util.cc b/libspu/kernel/test_util.cc
new file mode 100644
index 00000000..2ea31c68
--- /dev/null
+++ b/libspu/kernel/test_util.cc
@@ -0,0 +1,60 @@
+// Copyright 2021 Ant Group Co., Ltd.
+//
+// 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 "libspu/kernel/test_util.h"
+
+#include "libspu/core/config.h"
+#include "libspu/core/encoding.h"
+#include "libspu/kernel/hal/constants.h"  // bad reference
+#include "libspu/mpc/factory.h"
+
+namespace spu::kernel::test {
+
+SPUContext makeSPUContext(RuntimeConfig config,
+                          const std::shared_ptr<yacl::link::Context>& lctx) {
+  populateRuntimecConfig(config);
+
+  SPUContext ctx(config, lctx);
+  mpc::Factory::RegisterProtocol(&ctx, lctx);
+
+  return ctx;
+}
+
+SPUContext makeSPUContext(ProtocolKind prot_kind, FieldType field,
+                          const std::shared_ptr<yacl::link::Context>& lctx) {
+  RuntimeConfig cfg;
+  cfg.set_protocol(prot_kind);
+  cfg.set_field(field);
+  cfg.set_enable_action_trace(false);
+
+  return makeSPUContext(cfg, lctx);
+}
+
+Value makeValue(SPUContext* ctx, PtBufferView init, Visibility vtype,
+                DataType dtype, ShapeView shape) {
+  if (dtype == DT_INVALID) {
+    dtype = getEncodeType(init.pt_type);
+  }
+  auto res = hal::constant(ctx, init, dtype, shape);
+  switch (vtype) {
+    case VIS_PUBLIC:
+      return res;
+    case VIS_SECRET:
+      return hal::_p2s(ctx, res).setDtype(res.dtype());
+    default:
+      SPU_THROW("not supported vtype={}", vtype);
+  }
+}
+
+}  // namespace spu::kernel::test
diff --git a/libspu/kernel/hal/test_util.h b/libspu/kernel/test_util.h
similarity index 66%
rename from libspu/kernel/hal/test_util.h
rename to libspu/kernel/test_util.h
index 404c055c..1c809640 100644
--- a/libspu/kernel/hal/test_util.h
+++ b/libspu/kernel/test_util.h
@@ -14,18 +14,24 @@
 
 #include "xtensor/xrandom.hpp"
 
+#include "libspu/core/context.h"
+#include "libspu/core/shape_util.h"
+#include "libspu/core/value.h"
 #include "libspu/core/xt_helper.h"
-#include "libspu/kernel/hal/constants.h"
-#include "libspu/kernel/hal/prot_wrapper.h"
-#include "libspu/kernel/hal/public_helper.h"
+#include "libspu/kernel/hal/prot_wrapper.h"   // bad reference
+#include "libspu/kernel/hal/public_helper.h"  // bad reference
 
-namespace spu::kernel::hal::test {
+namespace spu::kernel::test {
 
-HalContext makeRefHalContext(RuntimeConfig config);
+SPUContext makeSPUContext(RuntimeConfig config,
+                          const std::shared_ptr<yacl::link::Context>& lctx);
 
-HalContext makeRefHalContext();
+SPUContext makeSPUContext(
+    ProtocolKind prot_kind = ProtocolKind::REF2K,
+    FieldType field = FieldType::FM64,
+    const std::shared_ptr<yacl::link::Context>& lctx = nullptr);
 
-Value makeValue(HalContext* ctx, PtBufferView init,
+Value makeValue(SPUContext* ctx, PtBufferView init,
                 Visibility vtype = VIS_PUBLIC, DataType dtype = DT_INVALID,
                 ShapeView shape = {});
 
@@ -42,16 +48,16 @@ auto xt_random(const std::vector<size_t>& shape, double min = -100,
   }
 }
 
-using UnaryOp = Value(HalContext*, const Value&);
-using BinaryOp = Value(HalContext*, const Value&, const Value&);
-using TernaryOp = Value(HalContext*, const Value&, const Value&, const Value&);
+using UnaryOp = Value(SPUContext*, const Value&);
+using BinaryOp = Value(SPUContext*, const Value&, const Value&);
+using TernaryOp = Value(SPUContext*, const Value&, const Value&, const Value&);
 
 template <typename T>
 xt::xarray<T> evalTernaryOp(Visibility in1_vtype, Visibility in2_vtype,
                             Visibility in3_vtype, TernaryOp* op,
                             PtBufferView in1, PtBufferView in2,
                             PtBufferView in3) {
-  HalContext ctx = makeRefHalContext();
+  SPUContext ctx = makeSPUContext();
 
   Value a = makeValue(&ctx, in1, in1_vtype);
   Value b = makeValue(&ctx, in2, in2_vtype);
@@ -60,17 +66,17 @@ xt::xarray<T> evalTernaryOp(Visibility in1_vtype, Visibility in2_vtype,
   Value d = op(&ctx, a, b, c);
 
   if (d.isSecret()) {
-    d = _s2p(&ctx, d);
+    d = hal::_s2p(&ctx, d);
   }
   SPU_ENFORCE(d.isPublic());
 
-  return dump_public_as<T>(&ctx, d);
+  return hal::dump_public_as<T>(&ctx, d);
 }
 
 template <typename T>
 xt::xarray<T> evalBinaryOp(Visibility lhs_vtype, Visibility rhs_vtype,
                            BinaryOp* op, PtBufferView lhs, PtBufferView rhs) {
-  HalContext ctx = makeRefHalContext();
+  SPUContext ctx = makeSPUContext();
 
   Value a = makeValue(&ctx, lhs, lhs_vtype);
   Value b = makeValue(&ctx, rhs, rhs_vtype);
@@ -78,27 +84,27 @@ xt::xarray<T> evalBinaryOp(Visibility lhs_vtype, Visibility rhs_vtype,
   Value c = op(&ctx, a, b);
 
   if (c.isSecret()) {
-    c = _s2p(&ctx, c).setDtype(c.dtype());
+    c = hal::_s2p(&ctx, c).setDtype(c.dtype());
   }
   SPU_ENFORCE(c.isPublic());
 
-  return dump_public_as<T>(&ctx, c);
+  return hal::dump_public_as<T>(&ctx, c);
 }
 
 template <typename T>
 xt::xarray<T> evalUnaryOp(Visibility in_vtype, UnaryOp* op, PtBufferView in) {
-  HalContext ctx = makeRefHalContext();
+  SPUContext ctx = makeSPUContext();
 
   Value a = makeValue(&ctx, in, in_vtype);
 
   Value b = op(&ctx, a);
 
   if (b.isSecret()) {
-    b = _s2p(&ctx, b).setDtype(b.dtype());
+    b = hal::_s2p(&ctx, b).setDtype(b.dtype());
   }
   SPU_ENFORCE(b.isPublic());
 
-  return dump_public_as<T>(&ctx, b);
+  return hal::dump_public_as<T>(&ctx, b);
 }
 
-}  // namespace spu::kernel::hal::test
+}  // namespace spu::kernel::test
diff --git a/libspu/mpc/common/ab_api_test.cc b/libspu/mpc/ab_api_test.cc
similarity index 59%
rename from libspu/mpc/common/ab_api_test.cc
rename to libspu/mpc/ab_api_test.cc
index 9edc4f9c..d1571e02 100644
--- a/libspu/mpc/common/ab_api_test.cc
+++ b/libspu/mpc/ab_api_test.cc
@@ -12,11 +12,12 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "libspu/mpc/common/ab_api_test.h"
+#include "libspu/mpc/ab_api_test.h"
 
 #include "libspu/core/prelude.h"
 #include "libspu/core/shape_util.h"
-#include "libspu/mpc/common/ab_api.h"
+#include "libspu/mpc/ab_api.h"
+#include "libspu/mpc/api.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/kernel.h"
 #include "libspu/mpc/utils/ring_ops.h"
@@ -25,9 +26,25 @@
 namespace spu::mpc::test {
 namespace {
 
-constexpr int64_t kNumel = 1000;
+Shape kShape = {30, 40};
 const std::vector<size_t> kShiftBits = {0, 1, 2, 31, 32, 33, 64, 1000};
 
+#define EXPECT_VALUE_EQ(X, Y)                         \
+  {                                                   \
+    EXPECT_EQ((X).shape(), (Y).shape());              \
+    auto [x_data, x_shape, x_dtype] = UnwrapValue(X); \
+    auto [y_data, y_shape, y_dtype] = UnwrapValue(Y); \
+    EXPECT_TRUE(ring_all_equal(x_data, y_data));      \
+  }
+
+#define EXPECT_VALUE_ALMOST_EQ(X, Y, ERR)             \
+  {                                                   \
+    EXPECT_EQ((X).shape(), (Y).shape());              \
+    auto [x_data, x_shape, x_dtype] = UnwrapValue(X); \
+    auto [y_data, y_shape, y_dtype] = UnwrapValue(Y); \
+    EXPECT_TRUE(ring_all_equal(x_data, y_data, ERR)); \
+  }
+
 bool verifyCost(Kernel* kernel, std::string_view name, const ce::Params& params,
                 const Communicator::Stats& cost, size_t repeated = 1) {
   if (kernel->kind() == Kernel::Kind::Dynamic) {
@@ -64,70 +81,73 @@ bool verifyCost(Kernel* kernel, std::string_view name, const ce::Params& params,
 }
 
 bool verifyCost(Kernel* kernel, std::string_view name, FieldType field,
-                size_t numel, size_t npc, const Communicator::Stats& cost) {
+                const Shape& shape, size_t npc,
+                const Communicator::Stats& cost) {
   ce::Params params = {{"K", SizeOf(field) * 8}, {"N", npc}};
-  return verifyCost(kernel, name, params, cost, numel /*repeated*/);
+  return verifyCost(kernel, name, params, cost, shape.numel() /*repeated*/);
 }
 
 }  // namespace
 
-#define TEST_ARITHMETIC_BINARY_OP_AA(OP)                                \
-  TEST_P(ArithmeticTest, OP##AA) {                                      \
-    const auto factory = std::get<0>(GetParam());                       \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                \
-    const size_t npc = std::get<2>(GetParam());                         \
-                                                                        \
-    utils::simulate(                                                    \
-        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {    \
-          auto obj = factory(conf, lctx);                               \
-                                                                        \
-          /* GIVEN */                                                   \
-          auto p0 = rand_p(obj.get(), kNumel);                          \
-          auto p1 = rand_p(obj.get(), kNumel);                          \
-                                                                        \
-          /* WHEN */                                                    \
-          auto a0 = p2a(obj.get(), p0);                                 \
-          auto a1 = p2a(obj.get(), p1);                                 \
-          auto prev = obj->getState<Communicator>()->getStats();        \
-          auto tmp = OP##_aa(obj.get(), a0, a1);                        \
-          auto cost = obj->getState<Communicator>()->getStats() - prev; \
-          auto re = a2p(obj.get(), tmp);                                \
-          auto rp = OP##_pp(obj.get(), p0, p1);                         \
-                                                                        \
-          /* THEN */                                                    \
-          EXPECT_TRUE(ring_all_equal(re, rp));                          \
-          EXPECT_TRUE(verifyCost(obj->getKernel(#OP "_aa"), #OP "_aa",  \
-                                 conf.field(), kNumel, npc, cost));     \
-        });                                                             \
+#define TEST_ARITHMETIC_BINARY_OP_AA(OP)                                       \
+  TEST_P(ArithmeticTest, OP##AA) {                                             \
+    const auto factory = std::get<0>(GetParam());                              \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                       \
+    const size_t npc = std::get<2>(GetParam());                                \
+                                                                               \
+    utils::simulate(                                                           \
+        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {           \
+          auto obj = factory(conf, lctx);                                      \
+                                                                               \
+          /* GIVEN */                                                          \
+          auto p0 = rand_p(obj.get(), kShape);                                 \
+          auto p1 = rand_p(obj.get(), kShape);                                 \
+                                                                               \
+          /* WHEN */                                                           \
+          auto a0 = p2a(obj.get(), p0);                                        \
+          auto a1 = p2a(obj.get(), p1);                                        \
+          auto prev = obj->prot()->getState<Communicator>()->getStats();       \
+          auto tmp = OP##_aa(obj.get(), a0, a1);                               \
+          auto cost =                                                          \
+              obj->prot()->getState<Communicator>()->getStats() - prev;        \
+          auto re = a2p(obj.get(), tmp);                                       \
+          auto rp = OP##_pp(obj.get(), p0, p1);                                \
+                                                                               \
+          /* THEN */                                                           \
+          EXPECT_VALUE_EQ(re, rp);                                             \
+          EXPECT_TRUE(verifyCost(obj->prot()->getKernel(#OP "_aa"), #OP "_aa", \
+                                 conf.field(), kShape, npc, cost));            \
+        });                                                                    \
   }
 
-#define TEST_ARITHMETIC_BINARY_OP_AP(OP)                                \
-  TEST_P(ArithmeticTest, OP##AP) {                                      \
-    const auto factory = std::get<0>(GetParam());                       \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                \
-    const size_t npc = std::get<2>(GetParam());                         \
-                                                                        \
-    utils::simulate(                                                    \
-        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {    \
-          auto obj = factory(conf, lctx);                               \
-                                                                        \
-          /* GIVEN */                                                   \
-          auto p0 = rand_p(obj.get(), kNumel);                          \
-          auto p1 = rand_p(obj.get(), kNumel);                          \
-                                                                        \
-          /* WHEN */                                                    \
-          auto a0 = p2a(obj.get(), p0);                                 \
-          auto prev = obj->getState<Communicator>()->getStats();        \
-          auto tmp = OP##_ap(obj.get(), a0, p1);                        \
-          auto cost = obj->getState<Communicator>()->getStats() - prev; \
-          auto re = a2p(obj.get(), tmp);                                \
-          auto rp = OP##_pp(obj.get(), p0, p1);                         \
-                                                                        \
-          /* THEN */                                                    \
-          EXPECT_TRUE(ring_all_equal(re, rp));                          \
-          EXPECT_TRUE(verifyCost(obj->getKernel(#OP "_ap"), #OP "_ap",  \
-                                 conf.field(), kNumel, npc, cost));     \
-        });                                                             \
+#define TEST_ARITHMETIC_BINARY_OP_AP(OP)                                       \
+  TEST_P(ArithmeticTest, OP##AP) {                                             \
+    const auto factory = std::get<0>(GetParam());                              \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                       \
+    const size_t npc = std::get<2>(GetParam());                                \
+                                                                               \
+    utils::simulate(                                                           \
+        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {           \
+          auto obj = factory(conf, lctx);                                      \
+                                                                               \
+          /* GIVEN */                                                          \
+          auto p0 = rand_p(obj.get(), kShape);                                 \
+          auto p1 = rand_p(obj.get(), kShape);                                 \
+                                                                               \
+          /* WHEN */                                                           \
+          auto a0 = p2a(obj.get(), p0);                                        \
+          auto prev = obj->prot()->getState<Communicator>()->getStats();       \
+          auto tmp = OP##_ap(obj.get(), a0, p1);                               \
+          auto cost =                                                          \
+              obj->prot()->getState<Communicator>()->getStats() - prev;        \
+          auto re = a2p(obj.get(), tmp);                                       \
+          auto rp = OP##_pp(obj.get(), p0, p1);                                \
+                                                                               \
+          /* THEN */                                                           \
+          EXPECT_VALUE_EQ(re, rp);                                             \
+          EXPECT_TRUE(verifyCost(obj->prot()->getKernel(#OP "_ap"), #OP "_ap", \
+                                 conf.field(), kShape, npc, cost));            \
+        });                                                                    \
   }
 
 #define TEST_ARITHMETIC_BINARY_OP(OP) \
@@ -142,38 +162,38 @@ TEST_P(ArithmeticTest, MulA1B) {
   const RuntimeConfig& conf = std::get<1>(GetParam());
   const size_t npc = std::get<2>(GetParam());
 
-  const std::vector<int64_t> shape{5, 5};
-
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
     auto obj = factory(conf, lctx);
 
     // MulA1B available for aby3 only for now.
-    if (!obj->hasKernel("mul_a1b")) {
+    if (!obj->prot()->hasKernel("mul_a1b")) {
       return;
     }
 
+    const size_t K = spu::SizeOf(conf.field()) * 8;
+
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), calcNumel(shape));
-    auto p1 = rand_p(obj.get(), calcNumel(shape));
-    ring_bitmask_(p1, 0, 1);
+    auto p0 = rand_p(obj.get(), kShape);
+    auto p1 = rand_p(obj.get(), kShape);
+    p1 = rshift_p(obj.get(), p1, K - 1);
     auto a0 = p2a(obj.get(), p0);
     auto a1 = p2b(obj.get(), p1);
     // hint runtime this is a 1bit value.
-    a1 = lshift_b(obj.get(), a1, spu::SizeOf(conf.field()) * 8 - 1);
-    a1 = rshift_b(obj.get(), a1, spu::SizeOf(conf.field()) * 8 - 1);
+    a1 = lshift_b(obj.get(), a1, K - 1);
+    a1 = rshift_b(obj.get(), a1, K - 1);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto tmp = mul_a1b(obj.get(), a0, a1);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     auto r_aa = a2p(obj.get(), tmp);
     auto r_pp = mul_pp(obj.get(), p0, p1);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_aa, r_pp));
-    EXPECT_TRUE(verifyCost(obj->getKernel("mul_a1b"), "mul_a1b", conf.field(),
-                           calcNumel(shape), npc, cost));
+    EXPECT_VALUE_EQ(r_aa, r_pp);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("mul_a1b"), "mul_a1b",
+                           conf.field(), kShape, npc, cost));
   });
 }
 
@@ -185,35 +205,35 @@ TEST_P(ArithmeticTest, MatMulAP) {
   const int64_t M = 3;
   const int64_t K = 4;
   const int64_t N = 3;
-  const std::vector<int64_t> shape_A{M, K};
-  const std::vector<int64_t> shape_B{K, N};
-  const std::vector<int64_t> shape_C{M, N};
+  const Shape shape_A = {M, K};
+  const Shape shape_B = {K, N};
+  const Shape shape_C = {M, N};
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), calcNumel(shape_A));
-    auto p1 = rand_p(obj.get(), calcNumel(shape_B));
+    auto p0 = rand_p(obj.get(), shape_A);
+    auto p1 = rand_p(obj.get(), shape_B);
     auto a0 = p2a(obj.get(), p0);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto tmp = mmul_ap(obj.get(), a0, p1, M, N, K);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     auto r_aa = a2p(obj.get(), tmp);
     auto r_pp = mmul_pp(obj.get(), p0, p1, M, N, K);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_aa, r_pp));
+    EXPECT_VALUE_EQ(r_aa, r_pp);
     ce::Params params = {{"K", SizeOf(conf.field()) * 8},
                          {"N", npc},
                          {"m", M},
                          {"n", N},
                          {"k", K}};
-    EXPECT_TRUE(
-        verifyCost(obj->getKernel("mmul_ap"), "mmul_ap", params, cost, 1));
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("mmul_ap"), "mmul_ap", params,
+                           cost, 1));
   });
 }
 
@@ -225,36 +245,36 @@ TEST_P(ArithmeticTest, MatMulAA) {
   const int64_t M = 3;
   const int64_t K = 4;
   const int64_t N = 3;
-  const std::vector<int64_t> shape_A{M, K};
-  const std::vector<int64_t> shape_B{K, N};
-  const std::vector<int64_t> shape_C{M, N};
+  const Shape shape_A = {M, K};
+  const Shape shape_B = {K, N};
+  const Shape shape_C = {M, N};
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), calcNumel(shape_A));
-    auto p1 = rand_p(obj.get(), calcNumel(shape_B));
+    auto p0 = rand_p(obj.get(), shape_A);
+    auto p1 = rand_p(obj.get(), shape_B);
     auto a0 = p2a(obj.get(), p0);
     auto a1 = p2a(obj.get(), p1);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto tmp = mmul_aa(obj.get(), a0, a1, M, N, K);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     auto r_aa = a2p(obj.get(), tmp);
     auto r_pp = mmul_pp(obj.get(), p0, p1, M, N, K);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_aa, r_pp));
+    EXPECT_VALUE_EQ(r_aa, r_pp);
     ce::Params params = {{"K", SizeOf(conf.field()) * 8},
                          {"N", npc},
                          {"m", M},
                          {"n", N},
                          {"k", K}};
-    EXPECT_TRUE(
-        verifyCost(obj->getKernel("mmul_aa"), "mmul_aa", params, cost, 1));
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("mmul_aa"), "mmul_aa", params,
+                           cost, 1));
   });
 }
 
@@ -267,21 +287,21 @@ TEST_P(ArithmeticTest, NotA) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
     auto a0 = p2a(obj.get(), p0);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto r_a = not_a(obj.get(), a0);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     auto r_p = a2p(obj.get(), r_a);
     auto r_pp = a2p(obj.get(), not_a(obj.get(), a0));
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_p, r_pp));
-    EXPECT_TRUE(verifyCost(obj->getKernel("not_a"), "not_a", conf.field(),
-                           kNumel, npc, cost));
+    EXPECT_VALUE_EQ(r_p, r_pp);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("not_a"), "not_a",
+                           conf.field(), kShape, npc, cost));
   });
 }
 
@@ -294,7 +314,7 @@ TEST_P(ArithmeticTest, LShiftA) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
     auto a0 = p2a(obj.get(), p0);
 
     for (auto bits : kShiftBits) {
@@ -303,16 +323,16 @@ TEST_P(ArithmeticTest, LShiftA) {
         continue;
       }
       /* WHEN */
-      auto prev = obj->getState<Communicator>()->getStats();
+      auto prev = obj->prot()->getState<Communicator>()->getStats();
       auto tmp = lshift_a(obj.get(), a0, bits);
-      auto cost = obj->getState<Communicator>()->getStats() - prev;
+      auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
       auto r_b = a2p(obj.get(), tmp);
       auto r_p = lshift_p(obj.get(), p0, bits);
 
       /* THEN */
-      EXPECT_TRUE(ring_all_equal(r_b, r_p));
-      EXPECT_TRUE(verifyCost(obj->getKernel("lshift_a"), "lshift_a",
-                             conf.field(), kNumel, npc, cost));
+      EXPECT_VALUE_EQ(r_b, r_p);
+      EXPECT_TRUE(verifyCost(obj->prot()->getKernel("lshift_a"), "lshift_a",
+                             conf.field(), kShape, npc, cost));
     }
   });
 }
@@ -322,17 +342,23 @@ TEST_P(ArithmeticTest, TruncA) {
   const RuntimeConfig& conf = std::get<1>(GetParam());
   const size_t npc = std::get<2>(GetParam());
 
-  ArrayRef p0_large =
-      ring_rand_range(conf.field(), kNumel, -(1 << 28), -(1 << 27));
-  ArrayRef p0_small = ring_rand_range(conf.field(), kNumel, 1, 10000);
+  // ArrayRef p0_large =
+  //     ring_rand_range(conf.field(), kShape, -(1 << 28), -(1 << 27));
+  // ArrayRef p0_small = ring_rand_range(conf.field(), kShape, 1, 10000);
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
     auto obj = factory(conf, lctx);
-    ArrayRef p0;
-    if (!static_cast<TruncAKernel*>(obj->getKernel("trunc_a"))->hasMsbError()) {
-      p0 = p0_large;
+    auto* kernel =
+        static_cast<TruncAKernel*>(obj->prot()->getKernel("trunc_a"));
+
+    auto p0 = rand_p(obj.get(), kShape);
+
+    if (!kernel->hasMsbError()) {
+      // trunc requires MSB to be zero.
+      p0 = arshift_p(obj.get(), p0, 1);
     } else {
-      p0 = p0_small;
+      // has msb error, only use lowest 10 bits.
+      p0 = arshift_p(obj.get(), p0, SizeOf(conf.field()) * 8 - 10);
     }
 
     /* GIVEN */
@@ -340,17 +366,17 @@ TEST_P(ArithmeticTest, TruncA) {
     auto a0 = p2a(obj.get(), p0);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto a1 = trunc_a(obj.get(), a0, bits);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     auto r_a = a2p(obj.get(), a1);
     auto r_p = arshift_p(obj.get(), p0, bits);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_a, r_p, npc));
-    EXPECT_TRUE(verifyCost(obj->getKernel("trunc_a"), "trunc_a", conf.field(),
-                           kNumel, npc, cost));
+    EXPECT_VALUE_ALMOST_EQ(r_a, r_p, npc);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("trunc_a"), "trunc_a",
+                           conf.field(), kShape, npc, cost));
   });
 }
 
@@ -363,18 +389,18 @@ TEST_P(ArithmeticTest, P2A) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto a0 = p2a(obj.get(), p0);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
     auto p1 = a2p(obj.get(), a0);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(p0, p1));
-    EXPECT_TRUE(verifyCost(obj->getKernel("p2a"), "p2a", conf.field(), kNumel,
-                           npc, cost));
+    EXPECT_VALUE_EQ(p0, p1);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("p2a"), "p2a", conf.field(),
+                           kShape, npc, cost));
   });
 }
 
@@ -387,78 +413,80 @@ TEST_P(ArithmeticTest, A2P) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
 
     /* WHEN */
     auto a0 = p2a(obj.get(), p0);
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto p1 = a2p(obj.get(), a0);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(p0, p1));
-    EXPECT_TRUE(verifyCost(obj->getKernel("a2p"), "a2p", conf.field(), kNumel,
-                           npc, cost));
+    EXPECT_VALUE_EQ(p0, p1);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("a2p"), "a2p", conf.field(),
+                           kShape, npc, cost));
   });
 }
 
-#define TEST_BOOLEAN_BINARY_OP_BB(OP)                                   \
-  TEST_P(BooleanTest, OP##BB) {                                         \
-    const auto factory = std::get<0>(GetParam());                       \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                \
-    const size_t npc = std::get<2>(GetParam());                         \
-                                                                        \
-    utils::simulate(                                                    \
-        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {    \
-          auto obj = factory(conf, lctx);                               \
-                                                                        \
-          /* GIVEN */                                                   \
-          auto p0 = rand_p(obj.get(), kNumel);                          \
-          auto p1 = rand_p(obj.get(), kNumel);                          \
-                                                                        \
-          /* WHEN */                                                    \
-          auto b0 = p2b(obj.get(), p0);                                 \
-          auto b1 = p2b(obj.get(), p1);                                 \
-          auto prev = obj->getState<Communicator>()->getStats();        \
-          auto tmp = OP##_bb(obj.get(), b0, b1);                        \
-          auto cost = obj->getState<Communicator>()->getStats() - prev; \
-          auto re = b2p(obj.get(), tmp);                                \
-          auto rp = OP##_pp(obj.get(), p0, p1);                         \
-                                                                        \
-          /* THEN */                                                    \
-          EXPECT_TRUE(ring_all_equal(re, rp));                          \
-          EXPECT_TRUE(verifyCost(obj->getKernel(#OP "_bb"), #OP "_bb",  \
-                                 conf.field(), kNumel, npc, cost));     \
-        });                                                             \
+#define TEST_BOOLEAN_BINARY_OP_BB(OP)                                          \
+  TEST_P(BooleanTest, OP##BB) {                                                \
+    const auto factory = std::get<0>(GetParam());                              \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                       \
+    const size_t npc = std::get<2>(GetParam());                                \
+                                                                               \
+    utils::simulate(                                                           \
+        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {           \
+          auto obj = factory(conf, lctx);                                      \
+                                                                               \
+          /* GIVEN */                                                          \
+          auto p0 = rand_p(obj.get(), kShape);                                 \
+          auto p1 = rand_p(obj.get(), kShape);                                 \
+                                                                               \
+          /* WHEN */                                                           \
+          auto b0 = p2b(obj.get(), p0);                                        \
+          auto b1 = p2b(obj.get(), p1);                                        \
+          auto prev = obj->prot()->getState<Communicator>()->getStats();       \
+          auto tmp = OP##_bb(obj.get(), b0, b1);                               \
+          auto cost =                                                          \
+              obj->prot()->getState<Communicator>()->getStats() - prev;        \
+          auto re = b2p(obj.get(), tmp);                                       \
+          auto rp = OP##_pp(obj.get(), p0, p1);                                \
+                                                                               \
+          /* THEN */                                                           \
+          EXPECT_VALUE_EQ(re, rp);                                             \
+          EXPECT_TRUE(verifyCost(obj->prot()->getKernel(#OP "_bb"), #OP "_bb", \
+                                 conf.field(), kShape, npc, cost));            \
+        });                                                                    \
   }
 
-#define TEST_BOOLEAN_BINARY_OP_BP(OP)                                   \
-  TEST_P(BooleanTest, OP##BP) {                                         \
-    const auto factory = std::get<0>(GetParam());                       \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                \
-    const size_t npc = std::get<2>(GetParam());                         \
-                                                                        \
-    utils::simulate(                                                    \
-        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {    \
-          auto obj = factory(conf, lctx);                               \
-                                                                        \
-          /* GIVEN */                                                   \
-          auto p0 = rand_p(obj.get(), kNumel);                          \
-          auto p1 = rand_p(obj.get(), kNumel);                          \
-                                                                        \
-          /* WHEN */                                                    \
-          auto b0 = p2b(obj.get(), p0);                                 \
-          auto prev = obj->getState<Communicator>()->getStats();        \
-          auto tmp = OP##_bp(obj.get(), b0, p1);                        \
-          auto cost = obj->getState<Communicator>()->getStats() - prev; \
-          auto re = b2p(obj.get(), tmp);                                \
-          auto rp = OP##_pp(obj.get(), p0, p1);                         \
-                                                                        \
-          /* THEN */                                                    \
-          EXPECT_TRUE(ring_all_equal(re, rp));                          \
-          EXPECT_TRUE(verifyCost(obj->getKernel(#OP "_bp"), #OP "_bp",  \
-                                 conf.field(), kNumel, npc, cost));     \
-        });                                                             \
+#define TEST_BOOLEAN_BINARY_OP_BP(OP)                                          \
+  TEST_P(BooleanTest, OP##BP) {                                                \
+    const auto factory = std::get<0>(GetParam());                              \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                       \
+    const size_t npc = std::get<2>(GetParam());                                \
+                                                                               \
+    utils::simulate(                                                           \
+        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {           \
+          auto obj = factory(conf, lctx);                                      \
+                                                                               \
+          /* GIVEN */                                                          \
+          auto p0 = rand_p(obj.get(), kShape);                                 \
+          auto p1 = rand_p(obj.get(), kShape);                                 \
+                                                                               \
+          /* WHEN */                                                           \
+          auto b0 = p2b(obj.get(), p0);                                        \
+          auto prev = obj->prot()->getState<Communicator>()->getStats();       \
+          auto tmp = OP##_bp(obj.get(), b0, p1);                               \
+          auto cost =                                                          \
+              obj->prot()->getState<Communicator>()->getStats() - prev;        \
+          auto re = b2p(obj.get(), tmp);                                       \
+          auto rp = OP##_pp(obj.get(), p0, p1);                                \
+                                                                               \
+          /* THEN */                                                           \
+          EXPECT_VALUE_EQ(re, rp);                                             \
+          EXPECT_TRUE(verifyCost(obj->prot()->getKernel(#OP "_bp"), #OP "_bp", \
+                                 conf.field(), kShape, npc, cost));            \
+        });                                                                    \
   }
 
 #define TEST_BOOLEAN_BINARY_OP(OP) \
@@ -468,37 +496,38 @@ TEST_P(ArithmeticTest, A2P) {
 TEST_BOOLEAN_BINARY_OP(and)
 TEST_BOOLEAN_BINARY_OP(xor)
 
-#define TEST_UNARY_OP_WITH_BIT_B(OP)                                      \
-  TEST_P(BooleanTest, OP##B) {                                            \
-    const auto factory = std::get<0>(GetParam());                         \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                  \
-    const size_t npc = std::get<2>(GetParam());                           \
-                                                                          \
-    utils::simulate(                                                      \
-        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {      \
-          auto obj = factory(conf, lctx);                                 \
-                                                                          \
-          /* GIVEN */                                                     \
-          auto p0 = rand_p(obj.get(), kNumel);                            \
-          auto b0 = p2b(obj.get(), p0);                                   \
-                                                                          \
-          for (auto bits : kShiftBits) {                                  \
-            if (bits >= p0.elsize() * 8) {                                \
-              continue;                                                   \
-            }                                                             \
-            /* WHEN */                                                    \
-            auto prev = obj->getState<Communicator>()->getStats();        \
-            auto tmp = OP##_b(obj.get(), b0, bits);                       \
-            auto cost = obj->getState<Communicator>()->getStats() - prev; \
-            auto r_b = b2p(obj.get(), tmp);                               \
-            auto r_p = OP##_p(obj.get(), p0, bits);                       \
-                                                                          \
-            /* THEN */                                                    \
-            EXPECT_TRUE(ring_all_equal(r_b, r_p));                        \
-            EXPECT_TRUE(verifyCost(obj->getKernel(#OP "_b"), #OP "_b",    \
-                                   conf.field(), kNumel, npc, cost));     \
-          }                                                               \
-        });                                                               \
+#define TEST_UNARY_OP_WITH_BIT_B(OP)                                           \
+  TEST_P(BooleanTest, OP##B) {                                                 \
+    const auto factory = std::get<0>(GetParam());                              \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                       \
+    const size_t npc = std::get<2>(GetParam());                                \
+                                                                               \
+    utils::simulate(                                                           \
+        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {           \
+          auto obj = factory(conf, lctx);                                      \
+                                                                               \
+          /* GIVEN */                                                          \
+          auto p0 = rand_p(obj.get(), kShape);                                 \
+          auto b0 = p2b(obj.get(), p0);                                        \
+                                                                               \
+          for (auto bits : kShiftBits) {                                       \
+            if (bits >= p0.elsize() * 8) {                                     \
+              continue;                                                        \
+            }                                                                  \
+            /* WHEN */                                                         \
+            auto prev = obj->prot()->getState<Communicator>()->getStats();     \
+            auto tmp = OP##_b(obj.get(), b0, bits);                            \
+            auto cost =                                                        \
+                obj->prot()->getState<Communicator>()->getStats() - prev;      \
+            auto r_b = b2p(obj.get(), tmp);                                    \
+            auto r_p = OP##_p(obj.get(), p0, bits);                            \
+                                                                               \
+            /* THEN */                                                         \
+            EXPECT_VALUE_EQ(r_b, r_p);                                         \
+            EXPECT_TRUE(verifyCost(obj->prot()->getKernel(#OP "_b"), #OP "_b", \
+                                   conf.field(), kShape, npc, cost));          \
+          }                                                                    \
+        });                                                                    \
   }
 
 TEST_UNARY_OP_WITH_BIT_B(lshift)
@@ -514,18 +543,18 @@ TEST_P(BooleanTest, P2B) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto b0 = p2b(obj.get(), p0);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
     auto p1 = b2p(obj.get(), b0);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(p0, p1));
-    EXPECT_TRUE(verifyCost(obj->getKernel("p2b"), "p2b", conf.field(), kNumel,
-                           npc, cost));
+    EXPECT_VALUE_EQ(p0, p1);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("p2b"), "p2b", conf.field(),
+                           kShape, npc, cost));
   });
 }
 
@@ -538,18 +567,18 @@ TEST_P(BooleanTest, B2P) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
 
     /* WHEN */
     auto b0 = p2b(obj.get(), p0);
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto p1 = b2p(obj.get(), b0);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(p0, p1));
-    EXPECT_TRUE(verifyCost(obj->getKernel("b2p"), "b2p", conf.field(), kNumel,
-                           npc, cost));
+    EXPECT_VALUE_EQ(p0, p1);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("b2p"), "b2p", conf.field(),
+                           kShape, npc, cost));
   });
 }
 
@@ -562,23 +591,23 @@ TEST_P(BooleanTest, BitrevB) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
 
     /* WHEN */
     auto b0 = p2b(obj.get(), p0);
 
     for (size_t i = 0; i < SizeOf(conf.field()); i++) {
       for (size_t j = i; j < SizeOf(conf.field()); j++) {
-        auto prev = obj->getState<Communicator>()->getStats();
+        auto prev = obj->prot()->getState<Communicator>()->getStats();
         auto b1 = bitrev_b(obj.get(), b0, i, j);
-        auto cost = obj->getState<Communicator>()->getStats() - prev;
+        auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
         auto p1 = b2p(obj.get(), b1);
         auto pp1 = bitrev_p(obj.get(), p0, i, j);
-        EXPECT_TRUE(ring_all_equal(p1, pp1));
+        EXPECT_VALUE_EQ(p1, pp1);
 
-        EXPECT_TRUE(verifyCost(obj->getKernel("bitrev_b"), "bitrev_b",
-                               conf.field(), kNumel, npc, cost));
+        EXPECT_TRUE(verifyCost(obj->prot()->getKernel("bitrev_b"), "bitrev_b",
+                               conf.field(), kShape, npc, cost));
       }
     }
   });
@@ -593,18 +622,18 @@ TEST_P(ConversionTest, A2B) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
     auto a0 = p2a(obj.get(), p0);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto b1 = a2b(obj.get(), a0);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     /* THEN */
-    EXPECT_TRUE(verifyCost(obj->getKernel("a2b"), "a2b", conf.field(), kNumel,
-                           npc, cost));
-    EXPECT_TRUE(ring_all_equal(p0, b2p(obj.get(), b1)));
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("a2b"), "a2b", conf.field(),
+                           kShape, npc, cost));
+    EXPECT_VALUE_EQ(p0, b2p(obj.get(), b1));
   });
 }
 
@@ -617,19 +646,19 @@ TEST_P(ConversionTest, B2A) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
     auto a0 = p2a(obj.get(), p0);
 
     /* WHEN */
     auto b1 = a2b(obj.get(), a0);
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto a1 = b2a(obj.get(), b1);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     /* THEN */
-    EXPECT_TRUE(verifyCost(obj->getKernel("b2a"), "b2a", conf.field(), kNumel,
-                           npc, cost));
-    EXPECT_TRUE(ring_all_equal(p0, a2p(obj.get(), a1)));
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("b2a"), "b2a", conf.field(),
+                           kShape, npc, cost));
+    EXPECT_VALUE_EQ(p0, a2p(obj.get(), a1));
   });
 }
 
@@ -641,24 +670,24 @@ TEST_P(ConversionTest, MSB) {
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
     auto obj = factory(conf, lctx);
 
-    if (!obj->hasKernel("msb_a2b")) {
+    if (!obj->prot()->hasKernel("msb_a2b")) {
       return;
     }
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
     auto a0 = p2a(obj.get(), p0);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto b1 = msb_a2b(obj.get(), a0);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     /* THEN */
-    EXPECT_TRUE(verifyCost(obj->getKernel("msb_a2b"), "msb_a2b", conf.field(),
-                           kNumel, npc, cost));
-    EXPECT_TRUE(ring_all_equal(ring_rshift(p0, SizeOf(conf.field()) * 8 - 1),
-                               b2p(obj.get(), b1)));
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("msb_a2b"), "msb_a2b",
+                           conf.field(), kShape, npc, cost));
+    EXPECT_VALUE_EQ(rshift_p(obj.get(), p0, SizeOf(conf.field()) * 8 - 1),
+                    b2p(obj.get(), b1));
   });
 }
 
diff --git a/libspu/mpc/common/ab_api_test.h b/libspu/mpc/ab_api_test.h
similarity index 92%
rename from libspu/mpc/common/ab_api_test.h
rename to libspu/mpc/ab_api_test.h
index 00cd9c4c..be473ca3 100644
--- a/libspu/mpc/common/ab_api_test.h
+++ b/libspu/mpc/ab_api_test.h
@@ -15,8 +15,7 @@
 #include "gtest/gtest.h"
 #include "yacl/link/link.h"
 
-#include "libspu/mpc/common/api_test_params.h"
-#include "libspu/mpc/object.h"
+#include "libspu/mpc/api_test_params.h"
 
 namespace spu::mpc::test {
 
diff --git a/libspu/mpc/aby3/protocol_test.cc b/libspu/mpc/aby3/protocol_test.cc
index 4a2783f5..79a9bb19 100644
--- a/libspu/mpc/aby3/protocol_test.cc
+++ b/libspu/mpc/aby3/protocol_test.cc
@@ -14,14 +14,15 @@
 
 #include "libspu/mpc/aby3/protocol.h"
 
+#include "libspu/mpc/ab_api_test.h"
 #include "libspu/mpc/api_test.h"
-#include "libspu/mpc/common/ab_api_test.h"
 
 namespace spu::mpc::test {
 namespace {
 
 RuntimeConfig makeConfig(FieldType field) {
   RuntimeConfig conf;
+  conf.set_protocol(ProtocolKind::ABY3);
   conf.set_field(field);
   return conf;
 }
diff --git a/libspu/mpc/api_test.cc b/libspu/mpc/api_test.cc
index ee963fce..6b1ee115 100644
--- a/libspu/mpc/api_test.cc
+++ b/libspu/mpc/api_test.cc
@@ -25,58 +25,74 @@
 namespace spu::mpc::test {
 namespace {
 
-constexpr int64_t kNumel = 1000;
+Shape kShape = {30, 40};
 const std::vector<size_t> kShiftBits = {0, 1, 2, 31, 32, 33, 64, 1000};
 
-}  // namespace
+#define EXPECT_VALUE_EQ(X, Y)                         \
+  {                                                   \
+    EXPECT_EQ((X).shape(), (Y).shape());              \
+    auto [x_data, x_shape, x_dtype] = UnwrapValue(X); \
+    auto [y_data, y_shape, y_dtype] = UnwrapValue(Y); \
+    EXPECT_TRUE(ring_all_equal(x_data, y_data));      \
+  }
 
-#define TEST_BINARY_OP_SS(OP)                                             \
-  TEST_P(ApiTest, OP##_ss) {                                              \
-    const auto factory = std::get<0>(GetParam());                         \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                  \
-    const size_t npc = std::get<2>(GetParam());                           \
-                                                                          \
-    utils::simulate(                                                      \
-        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {      \
-          auto obj = factory(conf, lctx);                                 \
-                                                                          \
-          /* GIVEN */                                                     \
-          auto p0 = rand_p(obj.get(), kNumel);                            \
-          auto p1 = rand_p(obj.get(), kNumel);                            \
-                                                                          \
-          /* WHEN */                                                      \
-          auto tmp =                                                      \
-              OP##_ss(obj.get(), p2s(obj.get(), p0), p2s(obj.get(), p1)); \
-          auto re = s2p(obj.get(), tmp);                                  \
-          auto rp = OP##_pp(obj.get(), p0, p1);                           \
-                                                                          \
-          /* THEN */                                                      \
-          EXPECT_TRUE(ring_all_equal(re, rp));                            \
-        });                                                               \
+#define EXPECT_VALUE_ALMOST_EQ(X, Y, ERR)             \
+  {                                                   \
+    EXPECT_EQ((X).shape(), (Y).shape());              \
+    auto [x_data, x_shape, x_dtype] = UnwrapValue(X); \
+    auto [y_data, y_shape, y_dtype] = UnwrapValue(Y); \
+    EXPECT_TRUE(ring_all_equal(x_data, y_data, ERR)); \
   }
 
-#define TEST_BINARY_OP_SP(OP)                                                \
-  TEST_P(ApiTest, OP##_sp) {                                                 \
+}  // namespace
+
+#define TEST_BINARY_OP_SS(OP)                                                \
+  TEST_P(ApiTest, OP##_ss) {                                                 \
     const auto factory = std::get<0>(GetParam());                            \
     const RuntimeConfig& conf = std::get<1>(GetParam());                     \
     const size_t npc = std::get<2>(GetParam());                              \
                                                                              \
-    utils::simulate(npc,                                                     \
-                    [&](const std::shared_ptr<yacl::link::Context>& lctx) {  \
-                      auto obj = factory(conf, lctx);                        \
+    utils::simulate(                                                         \
+        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {         \
+          auto sctx = factory(conf, lctx);                                   \
                                                                              \
-                      /* GIVEN */                                            \
-                      auto p0 = rand_p(obj.get(), kNumel);                   \
-                      auto p1 = rand_p(obj.get(), kNumel);                   \
+          /* GIVEN */                                                        \
+          auto p0 = rand_p(sctx.get(), kShape);                              \
+          auto p1 = rand_p(sctx.get(), kShape);                              \
                                                                              \
-                      /* WHEN */                                             \
-                      auto tmp = OP##_sp(obj.get(), p2s(obj.get(), p0), p1); \
-                      auto re = s2p(obj.get(), tmp);                         \
-                      auto rp = OP##_pp(obj.get(), p0, p1);                  \
+          /* WHEN */                                                         \
+          auto tmp =                                                         \
+              OP##_ss(sctx.get(), p2s(sctx.get(), p0), p2s(sctx.get(), p1)); \
+          auto re = s2p(sctx.get(), tmp);                                    \
+          auto rp = OP##_pp(sctx.get(), p0, p1);                             \
                                                                              \
-                      /* THEN */                                             \
-                      EXPECT_TRUE(ring_all_equal(re, rp));                   \
-                    });                                                      \
+          /* THEN */                                                         \
+          EXPECT_VALUE_EQ(re, rp);                                           \
+        });                                                                  \
+  }
+
+#define TEST_BINARY_OP_SP(OP)                                                  \
+  TEST_P(ApiTest, OP##_sp) {                                                   \
+    const auto factory = std::get<0>(GetParam());                              \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                       \
+    const size_t npc = std::get<2>(GetParam());                                \
+                                                                               \
+    utils::simulate(npc,                                                       \
+                    [&](const std::shared_ptr<yacl::link::Context>& lctx) {    \
+                      auto sctx = factory(conf, lctx);                         \
+                                                                               \
+                      /* GIVEN */                                              \
+                      auto p0 = rand_p(sctx.get(), kShape);                    \
+                      auto p1 = rand_p(sctx.get(), kShape);                    \
+                                                                               \
+                      /* WHEN */                                               \
+                      auto tmp = OP##_sp(sctx.get(), p2s(sctx.get(), p0), p1); \
+                      auto re = s2p(sctx.get(), tmp);                          \
+                      auto rp = OP##_pp(sctx.get(), p0, p1);                   \
+                                                                               \
+                      /* THEN */                                               \
+                      EXPECT_VALUE_EQ(re, rp);                                 \
+                    });                                                        \
   }
 
 #define TEST_BINARY_OP(OP) \
@@ -88,26 +104,26 @@ TEST_BINARY_OP(mul)
 TEST_BINARY_OP(and)
 TEST_BINARY_OP(xor)
 
-#define TEST_UNARY_OP_S(OP)                                                 \
-  TEST_P(ApiTest, OP##_s) {                                                 \
-    const auto factory = std::get<0>(GetParam());                           \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                    \
-    const size_t npc = std::get<2>(GetParam());                             \
-                                                                            \
-    utils::simulate(                                                        \
-        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {        \
-          auto obj = factory(conf, lctx);                                   \
-                                                                            \
-          /* GIVEN */                                                       \
-          auto p0 = rand_p(obj.get(), kNumel);                              \
-                                                                            \
-          /* WHEN */                                                        \
-          auto r_s = s2p(obj.get(), OP##_s(obj.get(), p2s(obj.get(), p0))); \
-          auto r_p = OP##_p(obj.get(), p0);                                 \
-                                                                            \
-          /* THEN */                                                        \
-          EXPECT_TRUE(ring_all_equal(r_s, r_p));                            \
-        });                                                                 \
+#define TEST_UNARY_OP_S(OP)                                                    \
+  TEST_P(ApiTest, OP##_s) {                                                    \
+    const auto factory = std::get<0>(GetParam());                              \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                       \
+    const size_t npc = std::get<2>(GetParam());                                \
+                                                                               \
+    utils::simulate(                                                           \
+        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {           \
+          auto sctx = factory(conf, lctx);                                     \
+                                                                               \
+          /* GIVEN */                                                          \
+          auto p0 = rand_p(sctx.get(), kShape);                                \
+                                                                               \
+          /* WHEN */                                                           \
+          auto r_s = s2p(sctx.get(), OP##_s(sctx.get(), p2s(sctx.get(), p0))); \
+          auto r_p = OP##_p(sctx.get(), p0);                                   \
+                                                                               \
+          /* THEN */                                                           \
+          EXPECT_VALUE_EQ(r_s, r_p);                                           \
+        });                                                                    \
   }
 
 #define TEST_UNARY_OP_P(OP)                                                 \
@@ -118,17 +134,17 @@ TEST_BINARY_OP(xor)
                                                                             \
     utils::simulate(npc,                                                    \
                     [&](const std::shared_ptr<yacl::link::Context>& lctx) { \
-                      auto obj = factory(conf, lctx);                       \
+                      auto sctx = factory(conf, lctx);                      \
                                                                             \
                       /* GIVEN */                                           \
-                      auto p0 = rand_p(obj.get(), kNumel);                  \
+                      auto p0 = rand_p(sctx.get(), kShape);                 \
                                                                             \
                       /* WHEN */                                            \
-                      auto r_p = OP##_p(obj.get(), p0);                     \
-                      auto r_pp = OP##_p(obj.get(), p0);                    \
+                      auto r_p = OP##_p(sctx.get(), p0);                    \
+                      auto r_pp = OP##_p(sctx.get(), p0);                   \
                                                                             \
                       /* THEN */                                            \
-                      EXPECT_TRUE(ring_all_equal(r_p, r_pp));               \
+                      EXPECT_VALUE_EQ(r_p, r_pp);                           \
                     });                                                     \
   }
 
@@ -139,32 +155,32 @@ TEST_BINARY_OP(xor)
 TEST_UNARY_OP(not )
 TEST_UNARY_OP(msb)
 
-#define TEST_UNARY_OP_WITH_BIT_S(OP)                                         \
-  TEST_P(ApiTest, OP##S) {                                                   \
-    const auto factory = std::get<0>(GetParam());                            \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                     \
-    const size_t npc = std::get<2>(GetParam());                              \
-                                                                             \
-    utils::simulate(                                                         \
-        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {         \
-          auto obj = factory(conf, lctx);                                    \
-                                                                             \
-          /* GIVEN */                                                        \
-          auto p0 = rand_p(obj.get(), kNumel);                               \
-                                                                             \
-          for (auto bits : kShiftBits) {                                     \
-            if (bits >= SizeOf(conf.field()) * 8) {                          \
-              continue;                                                      \
-            }                                                                \
-            /* WHEN */                                                       \
-            auto r_s =                                                       \
-                s2p(obj.get(), OP##_s(obj.get(), p2s(obj.get(), p0), bits)); \
-            auto r_p = OP##_p(obj.get(), p0, bits);                          \
-                                                                             \
-            /* THEN */                                                       \
-            EXPECT_TRUE(ring_all_equal(r_s, r_p));                           \
-          }                                                                  \
-        });                                                                  \
+#define TEST_UNARY_OP_WITH_BIT_S(OP)                                       \
+  TEST_P(ApiTest, OP##S) {                                                 \
+    const auto factory = std::get<0>(GetParam());                          \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                   \
+    const size_t npc = std::get<2>(GetParam());                            \
+                                                                           \
+    utils::simulate(                                                       \
+        npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {       \
+          auto sctx = factory(conf, lctx);                                 \
+                                                                           \
+          /* GIVEN */                                                      \
+          auto p0 = rand_p(sctx.get(), kShape);                            \
+                                                                           \
+          for (auto bits : kShiftBits) {                                   \
+            if (bits >= SizeOf(conf.field()) * 8) {                        \
+              continue;                                                    \
+            }                                                              \
+            /* WHEN */                                                     \
+            auto r_s = s2p(sctx.get(),                                     \
+                           OP##_s(sctx.get(), p2s(sctx.get(), p0), bits)); \
+            auto r_p = OP##_p(sctx.get(), p0, bits);                       \
+                                                                           \
+            /* THEN */                                                     \
+            EXPECT_VALUE_EQ(r_s, r_p);                                     \
+          }                                                                \
+        });                                                                \
   }
 
 #define TEST_UNARY_OP_WITH_BIT_P(OP)                                        \
@@ -175,20 +191,20 @@ TEST_UNARY_OP(msb)
                                                                             \
     utils::simulate(npc,                                                    \
                     [&](const std::shared_ptr<yacl::link::Context>& lctx) { \
-                      auto obj = factory(conf, lctx);                       \
+                      auto sctx = factory(conf, lctx);                      \
                                                                             \
                       /* GIVEN */                                           \
-                      auto p0 = rand_p(obj.get(), kNumel);                  \
+                      auto p0 = rand_p(sctx.get(), kShape);                 \
                                                                             \
                       for (auto bits : kShiftBits) { /* WHEN */             \
                         if (bits >= SizeOf(conf.field()) * 8) {             \
                           continue;                                         \
                         }                                                   \
-                        auto r_p = OP##_p(obj.get(), p0, bits);             \
-                        auto r_pp = OP##_p(obj.get(), p0, bits);            \
+                        auto r_p = OP##_p(sctx.get(), p0, bits);            \
+                        auto r_pp = OP##_p(sctx.get(), p0, bits);           \
                                                                             \
                         /* THEN */                                          \
-                        EXPECT_TRUE(ring_all_equal(r_p, r_pp));             \
+                        EXPECT_VALUE_EQ(r_p, r_pp);                         \
                       }                                                     \
                     });                                                     \
   }
@@ -207,16 +223,18 @@ TEST_P(ApiTest, TruncS) {
   const size_t npc = std::get<2>(GetParam());
 
   // trunc_pr only work for smalle range.
-  auto p0 = ring_rand_range(conf.field(), kNumel, 0, 10000);
+  auto p0_data = ring_rand_range(conf.field(), calcNumel(kShape), 0, 10000);
+  auto p0 = WrapValue(p0_data, kShape);
+
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
-    auto obj = factory(conf, lctx);
+    auto sctx = factory(conf, lctx);
 
     const size_t bits = 2;
-    auto r_s = s2p(obj.get(), trunc_s(obj.get(), p2s(obj.get(), p0), bits));
-    auto r_p = arshift_p(obj.get(), p0, bits);
+    auto r_s = s2p(sctx.get(), trunc_s(sctx.get(), p2s(sctx.get(), p0), bits));
+    auto r_p = arshift_p(sctx.get(), p0, bits);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_s, r_p, npc));
+    EXPECT_VALUE_ALMOST_EQ(r_s, r_p, npc);
   });
 }
 
@@ -229,33 +247,35 @@ TEST_P(ApiTest, MatMulSS) {
   const int64_t K = 400;
   const int64_t N = 60;
   const int64_t N2 = 90;
-  const std::vector<int64_t> shape_A{M, K};
-  const std::vector<int64_t> shape_B{K, N};
-  const std::vector<int64_t> shape_B2{K, N2};
+  const Shape shape_A = {M, K};
+  const Shape shape_B = {K, N};
+  const Shape shape_B2 = {K, N2};
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
-    auto obj = factory(conf, lctx);
+    auto sctx = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), calcNumel(shape_A));
-    auto p1 = rand_p(obj.get(), calcNumel(shape_B));
-    auto p2 = rand_p(obj.get(), calcNumel(shape_B2));
+    auto p0 = rand_p(sctx.get(), shape_A);
+    auto p1 = rand_p(sctx.get(), shape_B);
+    auto p2 = rand_p(sctx.get(), shape_B2);
+
+    auto s0 = p2s(sctx.get(), p0);
+    auto s1 = p2s(sctx.get(), p1);
+    auto s2 = p2s(sctx.get(), p2);
 
     /* WHEN */
-    auto tmp =
-        mmul_ss(obj.get(), p2s(obj.get(), p0), p2s(obj.get(), p1), M, N, K);
-    auto tmp2 =
-        mmul_ss(obj.get(), p2s(obj.get(), p0), p2s(obj.get(), p2), M, N2, K);
+    auto tmp = mmul_ss(sctx.get(), s0, s1, M, N, K);
+    auto tmp2 = mmul_ss(sctx.get(), s0, s2, M, N2, K);
 
-    auto r_ss = s2p(obj.get(), tmp);
-    auto r_pp = mmul_pp(obj.get(), p0, p1, M, N, K);
+    auto r_ss = s2p(sctx.get(), tmp);
+    auto r_pp = mmul_pp(sctx.get(), p0, p1, M, N, K);
 
-    auto r_ss2 = s2p(obj.get(), tmp2);
-    auto r_pp2 = mmul_pp(obj.get(), p0, p2, M, N2, K);
+    auto r_ss2 = s2p(sctx.get(), tmp2);
+    auto r_pp2 = mmul_pp(sctx.get(), p0, p2, M, N2, K);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_ss, r_pp));
-    EXPECT_TRUE(ring_all_equal(r_ss2, r_pp2));
+    EXPECT_VALUE_EQ(r_ss, r_pp);
+    EXPECT_VALUE_EQ(r_ss2, r_pp2);
   });
 }
 
@@ -267,23 +287,23 @@ TEST_P(ApiTest, MmulSP) {
   const int64_t M = 3;
   const int64_t K = 4;
   const int64_t N = 3;
-  const std::vector<int64_t> shape_A{M, K};
-  const std::vector<int64_t> shape_B{K, N};
+  const Shape shape_A = {M, K};
+  const Shape shape_B = {K, N};
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
-    auto obj = factory(conf, lctx);
+    auto sctx = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), calcNumel(shape_A));
-    auto p1 = rand_p(obj.get(), calcNumel(shape_B));
+    auto p0 = rand_p(sctx.get(), shape_A);
+    auto p1 = rand_p(sctx.get(), shape_B);
 
     /* WHEN */
-    auto tmp = mmul_sp(obj.get(), p2s(obj.get(), p0), p1, M, N, K);
-    auto r_ss = s2p(obj.get(), tmp);
-    auto r_pp = mmul_pp(obj.get(), p0, p1, M, N, K);
+    auto tmp = mmul_sp(sctx.get(), p2s(sctx.get(), p0), p1, M, N, K);
+    auto r_ss = s2p(sctx.get(), tmp);
+    auto r_pp = mmul_pp(sctx.get(), p0, p1, M, N, K);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_ss, r_pp));
+    EXPECT_VALUE_EQ(r_ss, r_pp);
   });
 }
 
@@ -293,17 +313,17 @@ TEST_P(ApiTest, P2S_S2P) {
   const size_t npc = std::get<2>(GetParam());
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
-    auto obj = factory(conf, lctx);
+    auto sctx = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(sctx.get(), kShape);
 
     /* WHEN */
-    auto s = p2s(obj.get(), p0);
-    auto p1 = s2p(obj.get(), s);
+    auto s = p2s(sctx.get(), p0);
+    auto p1 = s2p(sctx.get(), s);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(p0, p1));
+    EXPECT_VALUE_EQ(p0, p1);
   });
 }
 
diff --git a/libspu/mpc/api_test.h b/libspu/mpc/api_test.h
index c6ea3803..3b205d7c 100644
--- a/libspu/mpc/api_test.h
+++ b/libspu/mpc/api_test.h
@@ -17,8 +17,7 @@
 #include "gtest/gtest.h"
 #include "yacl/link/link.h"
 
-#include "libspu/mpc/common/api_test_params.h"
-#include "libspu/mpc/object.h"
+#include "libspu/mpc/api_test_params.h"
 
 #include "libspu/spu.pb.h"
 
diff --git a/libspu/mpc/common/api_test_params.h b/libspu/mpc/api_test_params.h
similarity index 90%
rename from libspu/mpc/common/api_test_params.h
rename to libspu/mpc/api_test_params.h
index b4b1d6b1..61277f44 100644
--- a/libspu/mpc/common/api_test_params.h
+++ b/libspu/mpc/api_test_params.h
@@ -16,7 +16,7 @@
 
 #include "yacl/link/link.h"
 
-#include "libspu/mpc/object.h"
+#include "libspu/core/context.h"
 
 namespace spu::mpc::test {
 
@@ -34,10 +34,10 @@ class NamedFunction : public std::function<Func> {
   const std::string& name() const { return name_; }
 };
 
-using CreateObjectFn = NamedFunction<std::unique_ptr<Object>(
+using CreateObjectFn = NamedFunction<std::unique_ptr<SPUContext>(
     const RuntimeConfig& conf,
     const std::shared_ptr<yacl::link::Context>& lctx)>;
 
 using OpTestParams = std::tuple<CreateObjectFn, RuntimeConfig, size_t>;
 
-}  // namespace spu::mpc::test
\ No newline at end of file
+}  // namespace spu::mpc::test
diff --git a/libspu/mpc/cheetah/protocol_test.cc b/libspu/mpc/cheetah/protocol_test.cc
index 6a3c9350..884c5ffb 100644
--- a/libspu/mpc/cheetah/protocol_test.cc
+++ b/libspu/mpc/cheetah/protocol_test.cc
@@ -14,14 +14,15 @@
 
 #include "libspu/mpc/cheetah/protocol.h"
 
+#include "libspu/mpc/ab_api_test.h"
 #include "libspu/mpc/api_test.h"
-#include "libspu/mpc/common/ab_api_test.h"
 
 namespace spu::mpc::test {
 namespace {
 
 RuntimeConfig makeConfig(FieldType field) {
   RuntimeConfig conf;
+  conf.set_protocol(ProtocolKind::CHEETAH);
   conf.set_field(field);
   return conf;
 }
diff --git a/libspu/mpc/ref2k/ref2k_test.cc b/libspu/mpc/ref2k/ref2k_test.cc
index dd2d5f15..19cef388 100644
--- a/libspu/mpc/ref2k/ref2k_test.cc
+++ b/libspu/mpc/ref2k/ref2k_test.cc
@@ -22,6 +22,7 @@ namespace {
 
 RuntimeConfig makeConfig(FieldType field) {
   RuntimeConfig conf;
+  conf.set_protocol(ProtocolKind::REF2K);
   conf.set_field(field);
   return conf;
 }
diff --git a/libspu/mpc/semi2k/protocol_test.cc b/libspu/mpc/semi2k/protocol_test.cc
index 48b78e67..439b1636 100644
--- a/libspu/mpc/semi2k/protocol_test.cc
+++ b/libspu/mpc/semi2k/protocol_test.cc
@@ -16,8 +16,8 @@
 
 #include <mutex>
 
+#include "libspu/mpc/ab_api_test.h"
 #include "libspu/mpc/api_test.h"
-#include "libspu/mpc/common/ab_api_test.h"
 #include "libspu/mpc/semi2k/beaver/ttp_server/beaver_server.h"
 
 namespace spu::mpc::test {
@@ -25,6 +25,7 @@ namespace {
 
 RuntimeConfig makeConfig(FieldType field) {
   RuntimeConfig conf;
+  conf.set_protocol(ProtocolKind::SEMI2K);
   conf.set_field(field);
   return conf;
 }
@@ -40,7 +41,7 @@ void InitBeaverServer() {
   });
 }
 
-std::unique_ptr<Object> makeTTPSemi2kProtocol(
+std::unique_ptr<SPUContext> makeTTPSemi2kProtocol(
     const RuntimeConfig& rt, const std::shared_ptr<yacl::link::Context>& lctx) {
   InitBeaverServer();
   RuntimeConfig ttp_rt = rt;
diff --git a/libspu/mpc/spdz2k/abprotocol_spdz2k_test.cc b/libspu/mpc/spdz2k/abprotocol_spdz2k_test.cc
index 7b8a0514..5af107d1 100644
--- a/libspu/mpc/spdz2k/abprotocol_spdz2k_test.cc
+++ b/libspu/mpc/spdz2k/abprotocol_spdz2k_test.cc
@@ -15,7 +15,8 @@
 #include "libspu/mpc/spdz2k/abprotocol_spdz2k_test.h"
 
 #include "libspu/core/shape_util.h"
-#include "libspu/mpc/common/ab_api.h"
+#include "libspu/mpc/ab_api.h"
+#include "libspu/mpc/api.h"
 #include "libspu/mpc/common/communicator.h"
 #include "libspu/mpc/kernel.h"
 #include "libspu/mpc/utils/ring_ops.h"
@@ -24,11 +25,30 @@
 namespace spu::mpc::test {
 namespace {
 
-constexpr int64_t kNumel = 10;
+Shape kShape = {30, 40};
 const std::vector<size_t> kShiftBits = {0, 1, 2, 31, 32, 33, 64, 1000};
 
+#define EXPECT_VALUE_EQ(X, Y)                         \
+  {                                                   \
+    EXPECT_EQ((X).dtype(), (Y).dtype());              \
+    EXPECT_EQ((X).shape(), (Y).shape());              \
+    auto [x_data, x_shape, x_dtype] = UnwrapValue(X); \
+    auto [y_data, y_shape, y_dtype] = UnwrapValue(Y); \
+    EXPECT_TRUE(ring_all_equal(x_data, y_data));      \
+  }
+
+#define EXPECT_VALUE_ALMOST_EQ(X, Y, ERR)             \
+  {                                                   \
+    EXPECT_EQ((X).dtype(), (Y).dtype());              \
+    EXPECT_EQ((X).shape(), (Y).shape());              \
+    auto [x_data, x_shape, x_dtype] = UnwrapValue(X); \
+    auto [y_data, y_shape, y_dtype] = UnwrapValue(Y); \
+    EXPECT_TRUE(ring_all_equal(x_data, y_data, ERR)); \
+  }
+
 bool verifyCost(Kernel* kernel, std::string_view name, FieldType field,
-                size_t numel, size_t npc, const Communicator::Stats& cost) {
+                const Shape& shape, size_t npc,
+                const Communicator::Stats& cost) {
   if (kernel->kind() == Kernel::Kind::Dynamic) {
     return true;
   }
@@ -36,6 +56,8 @@ bool verifyCost(Kernel* kernel, std::string_view name, FieldType field,
   auto comm = kernel->comm();
   auto latency = kernel->latency();
 
+  size_t numel = shape.numel();
+
   bool succeed = true;
   constexpr size_t kBitsPerBytes = 8;
   ce::Params params = {{"K", SizeOf(field) * 8}, {"N", npc}};
@@ -61,21 +83,21 @@ TEST_P(ArithmeticTest, P2A) {
   const size_t npc = std::get<2>(GetParam());
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
-    auto obj = factory(conf, lctx);
+    auto sctx = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(sctx.get(), kShape);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
-    auto a0 = p2a(obj.get(), p0);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
-    auto p1 = a2p(obj.get(), a0);
+    auto prev = sctx->prot()->getState<Communicator>()->getStats();
+    auto a0 = p2a(sctx.get(), p0);
+    auto cost = sctx->prot()->getState<Communicator>()->getStats() - prev;
+    auto p1 = a2p(sctx.get(), a0);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(p0, p1));
-    EXPECT_TRUE(verifyCost(obj->getKernel("p2a"), "p2a", conf.field(), kNumel,
-                           npc, cost));
+    EXPECT_VALUE_EQ(p0, p1);
+    EXPECT_TRUE(verifyCost(sctx->prot()->getKernel("p2a"), "p2a", conf.field(),
+                           kShape, npc, cost));
   });
 }
 
@@ -85,80 +107,80 @@ TEST_P(ArithmeticTest, A2P) {
   const size_t npc = std::get<2>(GetParam());
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
-    auto obj = factory(conf, lctx);
+    auto sctx = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(sctx.get(), kShape);
 
     /* WHEN */
-    auto a0 = p2a(obj.get(), p0);
-    auto prev = obj->getState<Communicator>()->getStats();
-    auto p1 = a2p(obj.get(), a0);
+    auto a0 = p2a(sctx.get(), p0);
+    auto prev = sctx->prot()->getState<Communicator>()->getStats();
+    auto p1 = a2p(sctx.get(), a0);
     [[maybe_unused]] auto cost =
-        obj->getState<Communicator>()->getStats() - prev;
+        sctx->prot()->getState<Communicator>()->getStats() - prev;
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(p0, p1));
-    EXPECT_TRUE(verifyCost(obj->getKernel("a2p"), "a2p", conf.field(), kNumel,
-                           npc, cost));
+    EXPECT_VALUE_EQ(p0, p1);
+    EXPECT_TRUE(verifyCost(sctx->prot()->getKernel("a2p"), "a2p", conf.field(),
+                           kShape, npc, cost));
   });
 }
 
-#define TEST_ARITHMETIC_BINARY_OP_AA(OP)                                  \
-  TEST_P(ArithmeticTest, OP##AA) {                                        \
-    const auto factory = std::get<0>(GetParam());                         \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                  \
-    const size_t npc = std::get<2>(GetParam());                           \
-                                                                          \
-    utils::simulate(npc, [&](std::shared_ptr<yacl::link::Context> lctx) { \
-      auto obj = factory(conf, lctx);                                     \
-                                                                          \
-      /* GIVEN */                                                         \
-      auto p0 = rand_p(obj.get(), kNumel);                                \
-      auto p1 = rand_p(obj.get(), kNumel);                                \
-                                                                          \
-      /* WHEN */                                                          \
-      auto a0 = p2a(obj.get(), p0);                                       \
-      auto a1 = p2a(obj.get(), p1);                                       \
-      auto prev = obj->getState<Communicator>()->getStats();              \
-      auto tmp = OP##_aa(obj.get(), a0, a1);                              \
-      auto cost = obj->getState<Communicator>()->getStats() - prev;       \
-      auto re = a2p(obj.get(), tmp);                                      \
-      auto rp = OP##_pp(obj.get(), p0, p1);                               \
-                                                                          \
-      /* THEN */                                                          \
-      EXPECT_TRUE(ring_all_equal(re, rp));                                \
-      EXPECT_TRUE(verifyCost(obj->getKernel(#OP "_aa"), #OP "_aa",        \
-                             conf.field(), kNumel, npc, cost));           \
-    });                                                                   \
+#define TEST_ARITHMETIC_BINARY_OP_AA(OP)                                    \
+  TEST_P(ArithmeticTest, OP##AA) {                                          \
+    const auto factory = std::get<0>(GetParam());                           \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                    \
+    const size_t npc = std::get<2>(GetParam());                             \
+                                                                            \
+    utils::simulate(npc, [&](std::shared_ptr<yacl::link::Context> lctx) {   \
+      auto obj = factory(conf, lctx);                                       \
+                                                                            \
+      /* GIVEN */                                                           \
+      auto p0 = rand_p(obj.get(), kShape);                                  \
+      auto p1 = rand_p(obj.get(), kShape);                                  \
+                                                                            \
+      /* WHEN */                                                            \
+      auto a0 = p2a(obj.get(), p0);                                         \
+      auto a1 = p2a(obj.get(), p1);                                         \
+      auto prev = obj->prot()->getState<Communicator>()->getStats();        \
+      auto tmp = OP##_aa(obj.get(), a0, a1);                                \
+      auto cost = obj->prot()->getState<Communicator>()->getStats() - prev; \
+      auto re = a2p(obj.get(), tmp);                                        \
+      auto rp = OP##_pp(obj.get(), p0, p1);                                 \
+                                                                            \
+      /* THEN */                                                            \
+      EXPECT_VALUE_EQ(re, rp);                                              \
+      EXPECT_TRUE(verifyCost(obj->prot()->getKernel(#OP "_aa"), #OP "_aa",  \
+                             conf.field(), kShape, npc, cost));             \
+    });                                                                     \
   }
 
-#define TEST_ARITHMETIC_BINARY_OP_AP(OP)                                  \
-  TEST_P(ArithmeticTest, OP##AP) {                                        \
-    const auto factory = std::get<0>(GetParam());                         \
-    const RuntimeConfig& conf = std::get<1>(GetParam());                  \
-    const size_t npc = std::get<2>(GetParam());                           \
-                                                                          \
-    utils::simulate(npc, [&](std::shared_ptr<yacl::link::Context> lctx) { \
-      auto obj = factory(conf, lctx);                                     \
-                                                                          \
-      /* GIVEN */                                                         \
-      auto p0 = rand_p(obj.get(), kNumel);                                \
-      auto p1 = rand_p(obj.get(), kNumel);                                \
-                                                                          \
-      /* WHEN */                                                          \
-      auto a0 = p2a(obj.get(), p0);                                       \
-      auto prev = obj->getState<Communicator>()->getStats();              \
-      auto tmp = OP##_ap(obj.get(), a0, p1);                              \
-      auto cost = obj->getState<Communicator>()->getStats() - prev;       \
-      auto re = a2p(obj.get(), tmp);                                      \
-      auto rp = OP##_pp(obj.get(), p0, p1);                               \
-                                                                          \
-      /* THEN */                                                          \
-      EXPECT_TRUE(ring_all_equal(re, rp));                                \
-      EXPECT_TRUE(verifyCost(obj->getKernel(#OP "_ap"), #OP "_ap",        \
-                             conf.field(), kNumel, npc, cost));           \
-    });                                                                   \
+#define TEST_ARITHMETIC_BINARY_OP_AP(OP)                                    \
+  TEST_P(ArithmeticTest, OP##AP) {                                          \
+    const auto factory = std::get<0>(GetParam());                           \
+    const RuntimeConfig& conf = std::get<1>(GetParam());                    \
+    const size_t npc = std::get<2>(GetParam());                             \
+                                                                            \
+    utils::simulate(npc, [&](std::shared_ptr<yacl::link::Context> lctx) {   \
+      auto obj = factory(conf, lctx);                                       \
+                                                                            \
+      /* GIVEN */                                                           \
+      auto p0 = rand_p(obj.get(), kShape);                                  \
+      auto p1 = rand_p(obj.get(), kShape);                                  \
+                                                                            \
+      /* WHEN */                                                            \
+      auto a0 = p2a(obj.get(), p0);                                         \
+      auto prev = obj->prot()->getState<Communicator>()->getStats();        \
+      auto tmp = OP##_ap(obj.get(), a0, p1);                                \
+      auto cost = obj->prot()->getState<Communicator>()->getStats() - prev; \
+      auto re = a2p(obj.get(), tmp);                                        \
+      auto rp = OP##_pp(obj.get(), p0, p1);                                 \
+                                                                            \
+      /* THEN */                                                            \
+      EXPECT_VALUE_EQ(re, rp);                                              \
+      EXPECT_TRUE(verifyCost(obj->prot()->getKernel(#OP "_ap"), #OP "_ap",  \
+                             conf.field(), kShape, npc, cost));             \
+    });                                                                     \
   }
 
 #define TEST_ARITHMETIC_BINARY_OP(OP) \
@@ -177,21 +199,21 @@ TEST_P(ArithmeticTest, NotA) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
     auto a0 = p2a(obj.get(), p0);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto r_a = not_a(obj.get(), a0);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     auto r_p = a2p(obj.get(), r_a);
     auto r_pp = a2p(obj.get(), not_a(obj.get(), a0));
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_p, r_pp));
-    EXPECT_TRUE(verifyCost(obj->getKernel("not_a"), "not_a", conf.field(),
-                           kNumel, npc, cost));
+    EXPECT_VALUE_EQ(r_p, r_pp);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("not_a"), "not_a",
+                           conf.field(), kShape, npc, cost));
   });
 }
 
@@ -203,30 +225,30 @@ TEST_P(ArithmeticTest, MatMulAP) {
   const int64_t M = 3;
   const int64_t K = 4;
   const int64_t N = 3;
-  const std::vector<int64_t> shape_A{M, K};
-  const std::vector<int64_t> shape_B{K, N};
-  const std::vector<int64_t> shape_C{M, N};
+  const Shape shape_A = {M, K};
+  const Shape shape_B = {K, N};
+  const Shape shape_C = {M, N};
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), calcNumel(shape_A));
-    auto p1 = rand_p(obj.get(), calcNumel(shape_B));
+    auto p0 = rand_p(obj.get(), shape_A);
+    auto p1 = rand_p(obj.get(), shape_B);
     auto a0 = p2a(obj.get(), p0);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto tmp = mmul_ap(obj.get(), a0, p1, M, N, K);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     auto r_aa = a2p(obj.get(), tmp);
     auto r_pp = mmul_pp(obj.get(), p0, p1, M, N, K);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_aa, r_pp));
-    EXPECT_TRUE(verifyCost(obj->getKernel("mmul_ap"), "mmul_ap", conf.field(),
-                           calcNumel(shape_C), npc, cost));
+    EXPECT_VALUE_EQ(r_aa, r_pp);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("mmul_ap"), "mmul_ap",
+                           conf.field(), shape_C, npc, cost));
   });
 }
 
@@ -238,31 +260,31 @@ TEST_P(ArithmeticTest, MatMulAA) {
   const int64_t M = 3;
   const int64_t K = 4;
   const int64_t N = 3;
-  const std::vector<int64_t> shape_A{M, K};
-  const std::vector<int64_t> shape_B{K, N};
-  const std::vector<int64_t> shape_C{M, N};
+  const Shape shape_A = {M, K};
+  const Shape shape_B = {K, N};
+  const Shape shape_C = {M, N};
 
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), calcNumel(shape_A));
-    auto p1 = rand_p(obj.get(), calcNumel(shape_B));
+    auto p0 = rand_p(obj.get(), shape_A);
+    auto p1 = rand_p(obj.get(), shape_B);
     auto a0 = p2a(obj.get(), p0);
     auto a1 = p2a(obj.get(), p1);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto tmp = mmul_aa(obj.get(), a0, a1, M, N, K);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     auto r_aa = a2p(obj.get(), tmp);
     auto r_pp = mmul_pp(obj.get(), p0, p1, M, N, K);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_aa, r_pp));
-    EXPECT_TRUE(verifyCost(obj->getKernel("mmul_aa"), "mmul_aa", conf.field(),
-                           calcNumel(shape_C), npc, cost));
+    EXPECT_VALUE_EQ(r_aa, r_pp);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("mmul_aa"), "mmul_aa",
+                           conf.field(), shape_C, npc, cost));
   });
 }
 
@@ -275,7 +297,7 @@ TEST_P(ArithmeticTest, LShiftA) {
     auto obj = factory(conf, lctx);
 
     /* GIVEN */
-    auto p0 = rand_p(obj.get(), kNumel);
+    auto p0 = rand_p(obj.get(), kShape);
     auto a0 = p2a(obj.get(), p0);
 
     for (auto bits : kShiftBits) {
@@ -284,16 +306,16 @@ TEST_P(ArithmeticTest, LShiftA) {
         continue;
       }
       /* WHEN */
-      auto prev = obj->getState<Communicator>()->getStats();
+      auto prev = obj->prot()->getState<Communicator>()->getStats();
       auto tmp = lshift_a(obj.get(), a0, bits);
-      auto cost = obj->getState<Communicator>()->getStats() - prev;
+      auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
       auto r_b = a2p(obj.get(), tmp);
       auto r_p = lshift_p(obj.get(), p0, bits);
 
       /* THEN */
-      EXPECT_TRUE(ring_all_equal(r_b, r_p));
-      EXPECT_TRUE(verifyCost(obj->getKernel("lshift_a"), "lshift_a",
-                             conf.field(), kNumel, npc, cost));
+      EXPECT_VALUE_EQ(r_b, r_p);
+      EXPECT_TRUE(verifyCost(obj->prot()->getKernel("lshift_a"), "lshift_a",
+                             conf.field(), kShape, npc, cost));
     }
   });
 }
@@ -303,35 +325,31 @@ TEST_P(ArithmeticTest, TruncA) {
   const RuntimeConfig& conf = std::get<1>(GetParam());
   const size_t npc = std::get<2>(GetParam());
 
-  ArrayRef p0_large =
-      ring_rand_range(conf.field(), kNumel, -(1 << 28), -(1 << 27));
-  ArrayRef p0_small = ring_rand_range(conf.field(), kNumel, 1, 10000);
-
   utils::simulate(npc, [&](const std::shared_ptr<yacl::link::Context>& lctx) {
     auto obj = factory(conf, lctx);
-    ArrayRef p0;
-    if (!static_cast<TruncAKernel*>(obj->getKernel("trunc_a"))->hasMsbError()) {
-      p0 = p0_large;
-    } else {
-      p0 = p0_small;
-    }
+    auto p0 = rand_p(obj.get(), kShape);
+
+    EXPECT_TRUE(static_cast<TruncAKernel*>(obj->prot()->getKernel("trunc_a"))
+                    ->hasMsbError());
+    // has msb error, only use lowest 10 bits.
+    p0 = arshift_p(obj.get(), p0, SizeOf(conf.field()) * 8 - 10);
 
     /* GIVEN */
     const size_t bits = 2;
     auto a0 = p2a(obj.get(), p0);
 
     /* WHEN */
-    auto prev = obj->getState<Communicator>()->getStats();
+    auto prev = obj->prot()->getState<Communicator>()->getStats();
     auto a1 = trunc_a(obj.get(), a0, bits);
-    auto cost = obj->getState<Communicator>()->getStats() - prev;
+    auto cost = obj->prot()->getState<Communicator>()->getStats() - prev;
 
     auto r_a = a2p(obj.get(), a1);
     auto r_p = arshift_p(obj.get(), p0, bits);
 
     /* THEN */
-    EXPECT_TRUE(ring_all_equal(r_a, r_p, npc));
-    EXPECT_TRUE(verifyCost(obj->getKernel("trunc_a"), "trunc_a", conf.field(),
-                           kNumel, npc, cost));
+    EXPECT_VALUE_ALMOST_EQ(r_a, r_p, npc);
+    EXPECT_TRUE(verifyCost(obj->prot()->getKernel("trunc_a"), "trunc_a",
+                           conf.field(), kShape, npc, cost));
   });
 }
 
diff --git a/libspu/mpc/spdz2k/abprotocol_spdz2k_test.h b/libspu/mpc/spdz2k/abprotocol_spdz2k_test.h
index 69a44f8b..b3d9698d 100644
--- a/libspu/mpc/spdz2k/abprotocol_spdz2k_test.h
+++ b/libspu/mpc/spdz2k/abprotocol_spdz2k_test.h
@@ -15,11 +15,11 @@
 #include "gtest/gtest.h"
 #include "yacl/link/link.h"
 
-#include "libspu/mpc/object.h"
+#include "libspu/core/context.h"
 
 namespace spu::mpc::test {
 
-using CreateObjectFn = std::function<std::unique_ptr<Object>(
+using CreateObjectFn = std::function<std::unique_ptr<SPUContext>(
     const RuntimeConfig& conf,
     const std::shared_ptr<yacl::link::Context>& lctx)>;
 
diff --git a/libspu/mpc/spdz2k/protocol_test.cc b/libspu/mpc/spdz2k/protocol_test.cc
index 6decec0c..96aaa2cd 100644
--- a/libspu/mpc/spdz2k/protocol_test.cc
+++ b/libspu/mpc/spdz2k/protocol_test.cc
@@ -21,6 +21,7 @@ namespace {
 
 RuntimeConfig makeConfig(FieldType field) {
   RuntimeConfig conf;
+  conf.set_protocol(ProtocolKind::SEMI2K);  // FIXME:
   conf.set_field(field);
   return conf;
 }
diff --git a/spu/tests/jnp_debug.py b/spu/tests/jnp_debug.py
index 0e7e6dde..b42d60e8 100644
--- a/spu/tests/jnp_debug.py
+++ b/spu/tests/jnp_debug.py
@@ -28,10 +28,13 @@
     """
 
     sim = ppsim.Simulator.simple(3, spu_pb2.ProtocolKind.ABY3, spu_pb2.FieldType.FM64)
+    copts = spu_pb2.CompilerOptions()
+    # Tweak compiler options
+    copts.disable_div_sqrt_rewrite = True
 
     x = np.random.randn(3, 4)
-    fn = lambda x: jnp.mean(x)
-    spu_fn = ppsim.sim_jax(sim, fn)
+    fn = lambda x: x / (jnp.sqrt(x) + 1e-9)
+    spu_fn = ppsim.sim_jax(sim, fn, copts=copts)
     z = spu_fn(x)
 
     print(spu_fn.pphlo)