Repository: secretflow/spu (language: python)

Issue:
spu进行函数计算，小数精度丢失异常
### Issue Type

Bug

### Source

binary

### Secretflow Version

0.8.2b2

### OS Platform and Distribution

Centos7 4c32g

### Python version

3.8.13

### Bazel version

_No response_

### GCC/Compiler version

_No response_

### What happend and What you expected to happen.

```shell
import json
import random
import re
import sys
import time

import secretflow as sf
import numpy as np
import pandas as pd
from dvadmin.dev_centre.sflib.item.init_data import init

def getFunAndElements(formulas: list):
    """
    将公式转为函数返回
    """
    map: dict = {}
    if formulas and len(formulas) > 0:
        for formula in formulas:
            formulaArray: list = formula.split("=")
            elements = re.findall(r'[a-zA-Z]+(?:_+[a-zA-Z]+)?\d*', formulaArray[1])
            map.update({formulaArray[0]: {"func": eval(f"(lambda {','.join(elements)}:{formulaArray[1]})"),
                                          "elements": elements}})
    return map


initdata = {
    "deviceList": [
        {
            "node_alias": "alice",
            "ip": "127.0.0.1",
            "port": "10003"
        },
        {
            "node_alias": "bob",
            "ip": "127.0.0.1",
            "port": "10004"
        }
    ],
    "address": "local",
    "selfDevice": {
        "node_alias": "alice",
        "ip": "127.0.0.1",
        "port": "10003"
    },
    "projectId": 000,
    "flowId": "flowId"
}
print(f"初始化数据：{initdata}")
spu, pyuMap, selfPyu = init(initdata)
for item in range(1, 2):
    # aa = [str(np.round(random.random()*100,2)) for i in range(1, 4)]
    # bb = [str(np.round(random.random()*100,2)) for i in range(1, 4)]
    aa=['199999.99','199999.99']
    bb=['0.00','100.01']
    df = pd.DataFrame({'A': aa, 'B': bb})
    print(f"df:{df}")
    #df.eval('C=A+B',inplace=True)
    funAndElement = getFunAndElements(['C=A+B','D=A-B'])
    funcC = funAndElement['C']["func"]
    funcD = funAndElement['D']["func"]
    data ={'A':pd.to_numeric(df['A']),'B':pd.to_numeric(df['B'])}
    spuObjC =spu(funcC)(**data)
    spuObjD =spu(funcD)(**data)

    item_result_C = sf.reveal(spuObjC).astype(str)
    item_result_D = sf.reveal(spuObjD).astype(str)

    print(f"item_result_C:{item_result_C}")
    print(f"item_result_D:{item_result_D}")
```


### Reproduction code to reproduce the issue.

```shell
df:           A       B
0  199999.99    0.00
1  199999.99  100.01

返回结果：
item_result_C:['199999.98' '200100.0']
item_result_D:['199999.98' '199899.97']

公式为：
C=A+B；  199999.99+0.00 却等于199999.98；
D=A-B， 199999.99-0.00 却等于199999.98； 199999.99-100.01却等于199899.97
```



Discussion / hints:
Hi @hy928302776 

spu 目前是用的 fixedpoint 来编码浮点数

0.99 是一个需要很多 fraction bits 才可以相对精确表达的一个值

如果有实际需求，可以尝试使用 fm128+比较多的 fraction bits 来一定程度的改善精度损失
这个有什么好的办法解决精度丢失问题吗？我这边该怎么做呢？上边说的 尝试使用 fm128+比较多的 fraction bits 来一定程度的改善精度损失  这个怎么做
> 这个有什么好的办法解决精度丢失问题吗？我这边该怎么做呢？上边说的 尝试使用 fm128+比较多的 fraction bits 来一定程度的改善精度损失 这个怎么做

用[这个](https://www.secretflow.org.cn/docs/secretflow/zh_CN/tutorial/spu_basics.html#%E8%BF%9B%E9%98%B6%E4%B8%BB%E9%A2%98%EF%BC%9A%E4%BD%BF%E7%94%A8%E4%B8%8D%E5%90%8C%E7%9A%84-MPC-%E5%8D%8F%E8%AE%AE)做个例子

```
    runtime_config={
        'protocol': spu.spu_pb2.ABY3,
        'field': spu.spu_pb2.FM128
    },
```

如果默认的依然不够精度，可以试一下加一个 

```
        'fxp_fraction_bits': 30
```

30 是一个可以配置的选项，可以根据自己的 fxp 需求计算一下大概需要多少 fxp bits
DEFAULT_SEMI2K_RUNTIME_CONFIG = {
    'protocol': spu.spu_pb2.SEMI2K,
    'field': spu.spu_pb2.FM128,
    'fxp_fraction_bits': 30,
}
我们的是两方计算，使用的这个协议配置，测试了下，还是不行，修改了fxp_fraction_bits的值也不行。计算：
 199999.97 - 100.01 结果返回还是 199899.95
> DEFAULT_SEMI2K_RUNTIME_CONFIG = { 'protocol': spu.spu_pb2.SEMI2K, 'field': spu.spu_pb2.FM128, 'fxp_fraction_bits': 30, } 我们的是两方计算，使用的这个协议配置，测试了下，还是不行，修改了fxp_fraction_bits的值也不行。计算： 199999.97 - 100.01 结果返回还是 199899.95

我仔细研究了一下现在大概是这样的的，用 FM64 举个例子：

199999.97 encode 到 fxp 是 52428791808
100.01 encode 到 fxp 是 26217022

52428791808 - 26217022 = 52402574786

result = 52402574786/(2^18)

这个数值在 f32 的时候是 199899.953125，所以结果出来会是 .95
如果用 f64, 这个值会是 199899.95874786376953

我们会在后续更新里尝试支持 f64 来解决这个精度问题


这个f64啥时间更新上去呢，我们着急使用发自我的 iPhone在 2023年5月18日，18:05，anakinxc ***@***.***> 写道：﻿

DEFAULT_SEMI2K_RUNTIME_CONFIG = { 'protocol': spu.spu_pb2.SEMI2K, 'field': spu.spu_pb2.FM128, 'fxp_fraction_bits': 30, } 我们的是两方计算，使用的这个协议配置，测试了下，还是不行，修改了fxp_fraction_bits的值也不行。计算： 199999.97 - 100.01 结果返回还是 199899.95

我仔细研究了一下现在大概是这样的的，用 FM64 举个例子：
199999.97 encode 到 fxp 是 52428791808
100.01 encode 到 fxp 是 26217022
52428791808 - 26217022 = 52402574786
result = 52402574786/(2^18)
这个数值在 f32 的时候是 199899.953125，所以结果出来会是 .95
如果用 f64, 这个值会是 199899.95874786376953
我们会在后续更新里尝试支持 f64 来解决这个精度问题

—Reply to this email directly, view it on GitHub, or unsubscribe.You are receiving this because you authored the thread.Message ID: ***@***.***>
> 这个f64啥时间更新上去呢，我们着急使用发自我的 iPhone在 2023年5月18日，18:05，anakinxc ***@***.***> 写道：﻿ DEFAULT_SEMI2K_RUNTIME_CONFIG = { 'protocol': spu.spu_pb2.SEMI2K, 'field': spu.spu_pb2.FM128, 'fxp_fraction_bits': 30, } 我们的是两方计算，使用的这个协议配置，测试了下，还是不行，修改了fxp_fraction_bits的值也不行。计算： 199999.97 - 100.01 结果返回还是 199899.95 我仔细研究了一下现在大概是这样的的，用 FM64 举个例子： 199999.97 encode 到 fxp 是 52428791808 100.01 encode 到 fxp 是 26217022 52428791808 - 26217022 = 52402574786 result = 52402574786/(2^18) 这个数值在 f32 的时候是 199899.953125，所以结果出来会是 .95 如果用 f64, 这个值会是 199899.95874786376953 我们会在后续更新里尝试支持 f64 来解决这个精度问题 —Reply to this email directly, view it on GitHub, or unsubscribe.You are receiving this because you authored the thread.Message ID: ***@***.***>

这个我们会尽快更新的

Fix patch:
diff --git a/.circleci/benchmark-config.yml b/.circleci/benchmark-config.yml
index db93b614..b02929b4 100644
--- a/.circleci/benchmark-config.yml
+++ b/.circleci/benchmark-config.yml
@@ -38,7 +38,7 @@ jobs:
       - run:
           name: Build SPU
           command: |
-            set +e 
+            set +e
             declare -i build_status
             sed -i "s/tensorflow /tensorflow-cpu /g" ./requirements-dev.txt
             docker run --name spu-build --mount type=bind,source="$(pwd)",target=/home/admin/dev/ secretflow/spu-ci:0.7 \
diff --git a/.circleci/benchmark.yml b/.circleci/benchmark.yml
index 7e5ebf62..a1d41e71 100644
--- a/.circleci/benchmark.yml
+++ b/.circleci/benchmark.yml
@@ -26,7 +26,7 @@ services:
       - -c
       - |
         cd /home/admin/dev
-        bazel-bin/examples/python/utils/nodectl -c .circleci/benchmark.json start --node_id node:0 
+        bazel-bin/examples/python/utils/nodectl -c .circleci/benchmark.json start --node_id node:0
     image: spu-build:v1
     networks:
       nn-benchmark:
@@ -39,7 +39,7 @@ services:
       - -c
       - |
         cd /home/admin/dev
-        bazel-bin/examples/python/utils/nodectl -c .circleci/benchmark.json start --node_id node:1 
+        bazel-bin/examples/python/utils/nodectl -c .circleci/benchmark.json start --node_id node:1
     image: spu-build:v1
     networks:
       nn-benchmark:
diff --git a/.circleci/continue-config.yml b/.circleci/continue-config.yml
index 73934d5d..cd0c703d 100644
--- a/.circleci/continue-config.yml
+++ b/.circleci/continue-config.yml
@@ -174,4 +174,3 @@ workflows:
     jobs:
       - linux_ut
       - macOS_ut
-      
\ No newline at end of file
diff --git a/.circleci/release-config.yml b/.circleci/release-config.yml
index 5bd01815..2635e834 100644
--- a/.circleci/release-config.yml
+++ b/.circleci/release-config.yml
@@ -116,12 +116,12 @@ jobs:
           command: |
             conda update -n base -c defaults conda -y
             conda create -n build python=<< parameters.python_ver >> -y
-            conda activate build            
+            conda activate build
 
             sh ./build_wheel_entrypoint.sh
             python3 -m pip install twine
             ls dist/*.whl
-            
+
             python3 -m twine upload -r pypi -u __token__ -p ${PYPI_TWINE_TOKEN} dist/*.whl
 
 # Invoke jobs via workflows
diff --git a/.github/workflows/cla.yml b/.github/workflows/cla.yml
index a162bb80..80eca44f 100644
--- a/.github/workflows/cla.yml
+++ b/.github/workflows/cla.yml
@@ -1,11 +1,11 @@
 # 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.
diff --git a/.github/workflows/trigger-circleci.yml b/.github/workflows/trigger-circleci.yml
index 422edc95..2bb70b2f 100644
--- a/.github/workflows/trigger-circleci.yml
+++ b/.github/workflows/trigger-circleci.yml
@@ -1,11 +1,11 @@
 # 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.
diff --git a/.markdownlint.yaml b/.markdownlint.yaml
new file mode 100644
index 00000000..ab9f1349
--- /dev/null
+++ b/.markdownlint.yaml
@@ -0,0 +1,59 @@
+# 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.
+
+# Enable all markdownlint rules by default
+default: true
+
+# List formatting:
+#
+# 1. Lists must have a blank line before AND after the list.
+# 2. Lists start aligned to the left (do not indent the top level list items).
+#    NOTE: markdownlint currently checks indentation for unordered lists only.
+#          Please manually verify that your ordered lists are not indented.
+#          See https://github.com/DavidAnson/markdownlint/issues/138.
+# 3. You may use one or zero blank lines between list items.
+# 4. Nested list items should be indented to align with the first character of
+#    the first line. For bullet lists, that means 2 spaces. For numbered
+#    lists, that's 3 spaces (but 4 spaces is okay if that's easier).
+# 5. In multiline list items, subsequent lines are indented by 2 spaces.
+#    This is not checked automatically, so we're documenting this convention
+#    to make sure the codebase stays consistent.
+#
+# Examples:
+#
+# * This is a list item that has multiple
+#   lines and each line aligns with the text from the first line.
+#   * This is a nested list, also aligned with the first line.
+#
+# For ordered lists, that means three spaces for wrapped lines:
+#
+# 1. This is an ordered list item.
+#    1. The nested list aligns with the first line.
+ul-indent:
+  indent: 2
+
+line-length:
+  line_length: 150
+  tables: false
+  code_blocks: false
+
+# Allow inline HTML
+no-inline-html: false
+
+# Allow dupe heading names only if they're not siblings
+no-duplicate-heading:
+  siblings_only: true
+
+# Allow images w/o alt-text
+no-alt-text: false
diff --git a/.vscode/settings.json b/.vscode/settings.json
index 52c78ac3..a026ff14 100644
--- a/.vscode/settings.json
+++ b/.vscode/settings.json
@@ -36,5 +36,6 @@
         "editor.formatOnSave": false,
     },
     "python.formatting.provider": "black",
-    "esbonio.sphinx.confDir": ""
+    "esbonio.sphinx.confDir": "",
+    "git.ignoreLimitWarning": true
 }
\ No newline at end of file
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 4c44fb8b..a4816fb0 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -7,7 +7,9 @@
 > - Add `[API]` prefix for API changes
 
 ## staging
+>
 > please add your unreleased change here.
+
 - [Feature] Improve secret dynamic slicing performance
 - [Feature] Improve accuracy of tanh approximation
 - [Example] Add flax GPT2
@@ -15,10 +17,11 @@
 - [Bugfix] Fix gen cache in unbalanced psi
 - [Bugfix] Fix possible compilation cache collision
 - [Improvement] Improve accuracy of mean with large size input
+- [Improvement] Improve support of Jax 64bit mode
 - [API] Add compiler options
 
-
 ## 20230406
+
 - [SPU] 0.3.2 release
 - [Feature] Add TrustedThirdParty beaver provider for semi2k
 - [Feature] Expose ssl/tls options
@@ -39,6 +42,7 @@
 - [PSI] Add ecdh-oprf 2-party two stage mode.
 
 ## 20230104
+
 - [SPU] 0.3.1 release
 - [API] Add get_var_meta
 - [API] Change ValueProto to bytes in Runtime/IO binding API
@@ -51,6 +55,7 @@
 - [3p] Update JAX to 0.4.1
 
 ## 20221116
+
 - [SPU] 0.3.0 release
 - [API] Add IR type for frontend compilation
 - [Feature] Support SignOp, Expm1Op, DotGeneralOp
@@ -62,11 +67,13 @@
 - [bugfix] Fix NotOp with non-pred integer types
 
 ## 20220905
+
 - [SPU] 0.2.4 release
 - [bugfix] Fix Tensorflow example
 - [miscs] Performance improvements
 
 ## 20220822
+
 - [SPU] 0.2.1 release
 - [Feature] Add compiler cache
 - [Feature] Support bc22-pcg psi
@@ -74,23 +81,26 @@
 - [miscs] Various performance improvements
 
 ## 20220727
+
 - [SPU] 0.1.1 release
 - [API] Added rsqrt support
 - [bugfix] Fixed some crashes and wrong answers
 
 ## 20220527
+
 - [SPU] 0.1.0 release
 - [API] Add exp/log approximation configurations.
 - [API] Change SigmoidMode enum value.
 
 ## 20220513
+
 - [SPU] 0.0.10
 - [Infra] Set minimum macOS requirements to 11.0
 - [Feature] Performance improvements
 - [Feature] Support SelectAndScatter.
 
-
 ## 20220511
+
 - [SPU] 0.0.9 release
 - [Feature] Add spu.binding.distributed.
 - [API] Expose visibility to ValueProto.
@@ -99,19 +109,23 @@
 - [bugfix] Ensure dynamic-slice & dynamic-update-slice OOB behavior is aligned with XLA semantics.
 
 ## 20220427
+
 - [SPU] 0.0.8.1 release
 - [bugfix] Fix iota type mismatch.
 
 ## 20220425
+
 - [SPU] 0.0.8 release
 - [Feature] Add tanh support.
 - [bugfix] Fix secret shift wrong answer.
 
 ## 20220424
+
 - [SPU] 0.0.7.1 release
 - [bugfix] Fix secret shift.
 
 ## 20220422
+
 - [SPU] 0.0.7 release
 - [Feature] Add python simulation module.
 - [Feature] Support general sort comparator.
@@ -119,33 +133,42 @@
 - [Feature] Support arithmetic right shift operation.
 - [API] Change DataType definition, support more accurate encoding.
 
-## 20220412
+## 20220412.1
+
 - [SPU] 0.0.6.1 release
 - [bugfix] Fixed several macOS build issues.
 
 ## 20220412
+
 - [SPU] 0.0.6 release
 - [API] Add storage_type/data_type, remove type_data.
-- [Experimental] Log(natural logarithm) is calculated with Pade approximation instead of Householder's approximation.
+- [Experimental] Log(natural logarithm) is calculated with Pade approximation
+                 instead of Householder's approximation.
 - [API] Fix typo, goldschmdit.
 - [Improvement] fxp div, improvement accuracy with normal distributed inputs.
 
 ## 20220325
+
 - [SPU] 0.0.5.1 release
 - [Bugfix] Fix SEMI2K divivsion wrong answer
 
 ## 20220324
+
 - [SPU] 0.0.5 release
 - [Improvement] Lift Convolution limitations.
 - [Improvement] Optimize maximum/minimum speed, with MPC B2A single bit conversion.
-- [Experimental] Div is calculated with direct goldschmdit method instead of multiplication of dividend and divisor reciprocal.
+- [Experimental] Div is calculated with direct goldschmdit method instead of
+                 multiplication of dividend and divisor reciprocal.
 - [API] change RuntimeConfig.enable_op_time_profile to enable_pphlo_profile.
 - [API] change RuntimeConfig.fxp_reciprocal_goldschmdit_iters to fxp_div_goldschmdit_iters.
 
 ## 20220308
+
 - [SPU] 0.0.4 release
-- [Feature] add silent ot support for various ot scenarios (chosen/correlated/random messages, chosen/correlated/random choices, 1o2/1oN)
-- [Feature] add non-linear computation protocols based on silent ot (comparison, truncation, b2a, triple, randbit, etc)
+- [Feature] add silent ot support for various ot scenarios (chosen/correlated/random
+            messages, chosen/correlated/random choices, 1o2/1oN)
+- [Feature] add non-linear computation protocols based on silent ot (comparison,
+            truncation, b2a, triple, randbit, etc)
 - [Feature] add a 2PC protocol: Cheetah
 - [Improvement] concatenate is a lot faster
 - [API] add RuntimeConfig.enable_op_time_profile to ask SPU to collect timing profiling data
@@ -162,19 +185,24 @@
 - [Feature] Left/right shift now properly supports non-scalar inputs
 
 ## 20220217
+
 - [spu] 0.0.2.3 release
 
 ## 20220211
+
 - [spu] 0.0.2.2 release
 - [Bugfix] Fix exception when kkrt psi input is too small
 
 ## 20220210
+
 - [spu] 0.0.2.1 release
 - [Bugfix] Fix exception when streaming psi output directory already exists
 
 ## 20220209
+
 - [spu] 0.0.2 release.
 - [Feature] Support multi-parties psi
 
 ## 20210930
+
 - [spu] Init release.
diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md
index 78226a3f..14a25013 100644
--- a/CONTRIBUTING.md
+++ b/CONTRIBUTING.md
@@ -17,28 +17,30 @@ part of the project.
 
 In general, please use clang-format to format code, and follow clang-tidy tips.
 
-Most of the code style is derived from the [Google C++ style guidelines](https://google.github.io/styleguide/cppguide.html), except:
+Most of the code style is derived from the
+[Google C++ style guidelines](https://google.github.io/styleguide/cppguide.html), except:
 
-* Exceptions are allowed and encouraged where appropriate.
-* Header guards should use `#pragma once`.
-* Adopt [camelBack](https://llvm.org/docs/Proposals/VariableNames.html#variable-names-coding-standard-options) for function names.
-* Use [fixed width integer types](https://en.cppreference.com/w/cpp/types/integer) whenever possible.
-* Avoid using size_t on interface APIs.
+- Exceptions are allowed and encouraged where appropriate.
+- Header guards should use `#pragma once`.
+- Adopt [camelBack](https://llvm.org/docs/Proposals/VariableNames.html#variable-names-coding-standard-options)
+    for function names.
+- Use [fixed width integer types](https://en.cppreference.com/w/cpp/types/integer) whenever possible.
+- Avoid using size_t on interface APIs.
 
 The compiler portion of the project follows [MLIR style](https://mlir.llvm.org/getting_started/DeveloperGuide/#style-guide).
 
 ### Other tips
 
-* Git commit message should be meaningful, we suggest imperative [keywords](https://github.com/joelparkerhenderson/git_commit_message#summary-keywords).
-* Developer must write unit-test (line coverage must be greater than 80%), tests should be deterministic.
-* Read awesome [Abseil Tips](https://abseil.io/tips/)
+- Git commit message should be meaningful, we suggest imperative [keywords](https://github.com/joelparkerhenderson/git_commit_message#summary-keywords).
+- Developer must write unit-test (line coverage must be greater than 80%), tests should be deterministic.
+- Read awesome [Abseil Tips](https://abseil.io/tips/)
 
 ## Build
 
 ### Prerequisite
 
-
 #### Docker
+
 ```sh
 ## start container
 docker run -d -it --name spu-gcc11-anolis-dev-$(whoami) \
@@ -90,6 +92,7 @@ pip install -r requirements-dev.txt
 ```
 
 ### Build & UnitTest
+
 ``` sh
 # build as debug
 bazel build //... -c dbg
@@ -114,15 +117,15 @@ bazel test //... --config=[macos-]ubsan
 
 ```sh
 # prerequisite
-pip install -U docs/requirements.txt
+pip install -U -r docs/requirements.txt
 
-cd docs & make html  # html docs will be in docs/_build/html
+cd docs && make html  # html docs will be in docs/_build/html
 ```
 
-
 ## Release cycle
 
-SPU recommends users "live-at-head" like [abseil-cpp](https://github.com/abseil/abseil-cpp), just like abseil, spu also provide Long Term Support Releases to which we backport fixes for severe bugs.
+SPU recommends users "live-at-head" like [abseil-cpp](https://github.com/abseil/abseil-cpp),
+just like abseil, spu also provide Long Term Support Releases to which we backport fixes for severe bugs.
 
 We use the release date as the version number, see [change log](CHANGELOG.md) for example.
 
@@ -130,5 +133,6 @@ We use the release date as the version number, see [change log](CHANGELOG.md) fo
 
 Please keep updating changes to the staging area of [change log](CHANGELOG.md)
 Changelog should contain:
+
 - all public API changes, including new features, breaking changes and deprecations.
 - notable internal changes, like performance improvement.
diff --git a/INSTALLATION.md b/INSTALLATION.md
index 10ae2048..a139452a 100644
--- a/INSTALLATION.md
+++ b/INSTALLATION.md
@@ -12,13 +12,13 @@ SPU has been tested with the following settings:
 
 ### MacOS
 
-We have conducted some successful preliminary testings on macOS Monterey 12.4 with Intel processors and Apple Silicon.
+We have conducted some successful preliminary testings on
+macOS Monterey 12.4 with Intel processors and Apple Silicon.
 
 ### Docker Image
 
 Please check [official Docker image](https://registry.hub.docker.com/r/secretflow/secretflow-gcc11-anolis-dev).
 
-
 ## Binaries
 
 ### From PyPI
@@ -29,7 +29,6 @@ You could install SPU via the [official PyPI package](https://pypi.org/project/s
 pip install spu
 ```
 
-
 ### From Source
 
 At the root of repo, run
diff --git a/LEGAL.md b/LEGAL.md
index f9689208..31408a13 100644
--- a/LEGAL.md
+++ b/LEGAL.md
@@ -1,7 +1,10 @@
-Legal Disclaimer
+# Legal Disclaimer
 
-Within this source code, the comments in Chinese shall be the original, governing version. Any comment in other languages are for reference only. In the event of any conflict between the Chinese language version comments and other language version comments, the Chinese language version shall prevail.
+Within this source code, the comments in Chinese shall be the original, governing version.
+Any comment in other languages are for reference only.
+In the event of any conflict between the Chinese language version comments and other
+language version comments, the Chinese language version shall prevail.
 
 法律免责声明
 
-关于代码注释部分，中文注释为官方版本，其它语言注释仅做参考。中文注释可能与其它语言注释存在不一致，当中文注释与其它语言注释存在不一致时，请以中文注释为准。
\ No newline at end of file
+关于代码注释部分，中文注释为官方版本，其它语言注释仅做参考。中文注释可能与其它语言注释存在不一致，当中文注释与其它语言注释存在不一致时，请以中文注释为准。
diff --git a/README.md b/README.md
index 942f2187..3f3265b1 100644
--- a/README.md
+++ b/README.md
@@ -2,13 +2,20 @@
 
 [![CircleCI](https://dl.circleci.com/status-badge/img/gh/secretflow/spu/tree/main.svg?style=svg)](https://dl.circleci.com/status-badge/redirect/gh/secretflow/spu/tree/main)
 
-SPU (Secure Processing Unit) aims to be a `provable`, `measurable` secure computation device, which provides computation ability while keeping your private data protected.
+SPU (Secure Processing Unit) aims to be a `provable`, `measurable` secure computation device,
+which provides computation ability while keeping your private data protected.
 
-SPU could be treated as a programmable device, it's not designed to be used directly. Normally we use SecretFlow framework, which use SPU as the underline secure computing device.
+SPU could be treated as a programmable device, it's not designed to be used directly.
+Normally we use SecretFlow framework, which use SPU as the underline secure computing device.
 
-Currently, we mainly focus on `provable` security. It contains a secure runtime that evaluates [XLA](https://www.tensorflow.org/xla/operation_semantics)-like tensor operations, which use [MPC](https://en.wikipedia.org/wiki/Secure_multi-party_computation) as the underline evaluation engine to protect privacy information.
+Currently, we mainly focus on `provable` security. It contains a secure runtime that evaluates
+[XLA](https://www.tensorflow.org/xla/operation_semantics)-like tensor operations,
+which use [MPC](https://en.wikipedia.org/wiki/Secure_multi-party_computation) as the underline
+evaluation engine to protect privacy information.
 
-SPU python package also contains a simple distributed module to demo SPU usage, but it's **NOT designed for production** due to system security and performance concerns, please **DO NOT** use it directly in production.
+SPU python package also contains a simple distributed module to demo SPU usage,
+but it's **NOT designed for production** due to system security and performance concerns,
+please **DO NOT** use it directly in production.
 
 ## Contribution Guidelines
 
diff --git a/REPO_LAYOUT.md b/REPO_LAYOUT.md
index 2dd77dee..4ea24f3c 100644
--- a/REPO_LAYOUT.md
+++ b/REPO_LAYOUT.md
@@ -8,18 +8,25 @@ This is a high level overview of how the repository is laid out. Some major fold
 * [docs/](docs/): Documents of SPU.
 * [examples/](examples/): Contains both cpp and python examples of SPU.
 * [libspu/](libspu/): Core C++ implementations of SPU.
-    - [compiler/](libspu/compiler/): SPU's compiler stack. It accepts the standard XLA IR along with inputs mpc metadata and lowering the XLA IR to an mpc specific IR.
-    - [core/](libspu/core/): Basic data structures used in SPU.
-    - [device/](libspu/device/): SPU device `Runtime`. It consists of IO(infeed, outfeed), symbol tables(storage) and the IR executor.
-    - [dialect/](libspu/dialect/): Internal MPC-specific IR used by SPU.
-    - [kernel/](libspu/kernel/): SPU runtime kernels.
-        * [hlo/](libspu/kernel/hlo/): High level operations, which match the definition of pphlo.
-        * [hal/](libspu/kernel/hal/): Hardware abstraction layer which implements crypto-independent core logics, fixed-point related abstractions and some non-linear APIs. It could be viewed as a builtin library in addition to the SPU VM.
-    - [mpc/](libspu/mpc/): Various mpc protocols. This folder defines the [standard interface](libspu/mpc/apis.h) different mpc protocols need to conform.
-        * [aby3/](libspu/mpc/aby3/): The semi-honest variant of ABY3 protocol. Currently only `Arithmetic` and `Boolean` are implemented.
-        * [cheetah/](libspu/mpc/cheetah/): An excellent semi-honest 2PC protocol implemented by [Alibaba-Gemini-Lab](https://alibaba-gemini-lab.github.io/).
-        * [semi2k/](libspu/mpc/semi2k/): The semi-honest variant of SPDZ protocol. It could be shipped with different correlated random generators.
-        * [ref2k/](libspu/mpc/ref2k/): A plaintext protocol. It is aimed to serve as a reference implementation of how a new protocol could be added in SPU.
-        * [util/](libspu/mpc/utils/): Common utilities for different mpc protocols.
+  * [compiler/](libspu/compiler/): SPU's compiler stack. It accepts the standard XLA IR along with
+                                   inputs mpc metadata and lowering the XLA IR to an mpc specific IR.
+  * [core/](libspu/core/): Basic data structures used in SPU.
+  * [device/](libspu/device/): SPU device `Runtime`. It consists of IO(infeed, outfeed),
+                               symbol tables(storage) and the IR executor.
+  * [dialect/](libspu/dialect/): Internal MPC-specific IR used by SPU.
+  * [kernel/](libspu/kernel/): SPU runtime kernels.
+    * [hlo/](libspu/kernel/hlo/): High level operations, which match the definition of pphlo.
+    * [hal/](libspu/kernel/hal/): Hardware abstraction layer which implements crypto-independent core
+                                  logics, fixed-point related abstractions and some non-linear APIs.
+                                  It could be viewed as a builtin library in addition to the SPU VM.
+  * [mpc/](libspu/mpc/): Various mpc protocols. This folder defines the [standard interface](libspu/mpc/apis.h)
+                         different mpc protocols need to conform.
+    * [aby3/](libspu/mpc/aby3/): The semi-honest variant of ABY3 protocol.
+                                 Currently only `Arithmetic` and `Boolean` are implemented.
+    * [cheetah/](libspu/mpc/cheetah/): An excellent semi-honest 2PC protocol implemented by [Alibaba-Gemini-Lab](https://alibaba-gemini-lab.github.io/).
+    * [semi2k/](libspu/mpc/semi2k/): The semi-honest variant of SPDZ protocol.
+                                     It could be shipped with different correlated random generators.
+    * [ref2k/](libspu/mpc/ref2k/): A plaintext protocol. It is aimed to serve as a reference implementation
+                                   of how a new protocol could be added in SPU.
+    * [util/](libspu/mpc/utils/): Common utilities for different mpc protocols.
 * [spu/](spu/): Python binding of libspu.
-
diff --git a/bazel/patches/apsi.patch b/bazel/patches/apsi.patch
index dc2a9802..c464acee 100644
--- a/bazel/patches/apsi.patch
+++ b/bazel/patches/apsi.patch
@@ -2,15 +2,6 @@ diff --git a/CMakeLists.txt b/CMakeLists.txt
 index 78d54a6..07df321 100644
 --- a/CMakeLists.txt
 +++ b/CMakeLists.txt
-@@ -122,7 +122,7 @@ set(APSI_INCLUDES_INSTALL_DIR ${CMAKE_INSTALL_INCLUDEDIR}/APSI-${APSI_VERSION_MA
- # find_package might throw a FATAL_ERROR before"xxx: not found", e.g. with vcpkg.cmake.
- 
- # Microsoft SEAL
--find_package(SEAL 4.1 QUIET REQUIRED)
-+find_package(SEAL 4.0 QUIET REQUIRED)
- if(NOT SEAL_FOUND)
-     message(FATAL_ERROR "Microsoft SEAL: not found")
- else()
 @@ -140,24 +140,6 @@ else()
      message(STATUS "Microsoft Kuku: found")
  endif()
diff --git a/bazel/repositories.bzl b/bazel/repositories.bzl
index e934bed4..b80100b5 100644
--- a/bazel/repositories.bzl
+++ b/bazel/repositories.bzl
@@ -138,8 +138,8 @@ def _com_github_xtensor_xtl():
     )
 
 def _com_github_openxla_xla():
-    OPENXLA_COMMIT = "6d294f24babbadbe7a42fbd91e9bead178acaa10"
-    OPENXLA_SHA256 = "9ead425b1159d32c42b4f5c3cbf48b9ec6331a288ceab6a034e419d0204a6229"
+    OPENXLA_COMMIT = "fbb42427fbc5acdb838ef66680b38ed5e2391211"
+    OPENXLA_SHA256 = "5b711639d527325e2b5068792af2862a0746adbab4a0c6291d73348402f242f8"
 
     SKYLIB_VERSION = "1.3.0"
 
@@ -193,11 +193,11 @@ def _com_intel_hexl():
         http_archive,
         name = "com_intel_hexl",
         type = "tar.gz",
-        strip_prefix = "hexl-1.2.3",
-        sha256 = "f2cf33ee2035d12996d10b69d2f41a586b9954a29b99c70a852495cf5758878c",
+        strip_prefix = "hexl-1.2.4",
+        sha256 = "f91dd158e3faa8795035f618286cdef7dacbf488fd0ba4f44409880997b4549f",
         build_file = "@spulib//bazel:hexl.BUILD",
         urls = [
-            "https://github.com/intel/hexl/archive/refs/tags/v1.2.3.tar.gz",
+            "https://github.com/intel/hexl/archive/refs/tags/v1.2.4.tar.gz",
         ],
     )
 
@@ -254,13 +254,13 @@ def _com_github_microsoft_seal():
     maybe(
         http_archive,
         name = "com_github_microsoft_seal",
-        sha256 = "616653498ba8f3e0cd23abef1d451c6e161a63bd88922f43de4b3595348b5c7e",
-        strip_prefix = "SEAL-4.0.0",
+        sha256 = "af9bf0f0daccda2a8b7f344f13a5692e0ee6a45fea88478b2b90c35648bf2672",
+        strip_prefix = "SEAL-4.1.1",
         type = "tar.gz",
         patch_args = ["-p1"],
         patches = ["@spulib//bazel:patches/seal.patch"],
         urls = [
-            "https://github.com/microsoft/SEAL/archive/refs/tags/v4.0.0.tar.gz",
+            "https://github.com/microsoft/SEAL/archive/refs/tags/v4.1.1.tar.gz",
         ],
         build_file = "@spulib//bazel:seal.BUILD",
     )
diff --git a/bazel/seal.BUILD b/bazel/seal.BUILD
index 98e972c3..e3098e42 100644
--- a/bazel/seal.BUILD
+++ b/bazel/seal.BUILD
@@ -57,8 +57,8 @@ spu_cmake_external(
         "//conditions:default": default_config,
     }),
     lib_source = "@com_github_microsoft_seal//:all",
-    out_include_dir = "include/SEAL-4.0",
-    out_static_libs = ["libseal-4.0.a"],
+    out_include_dir = "include/SEAL-4.1",
+    out_static_libs = ["libseal-4.1.a"],
     deps = [
         "@com_github_facebook_zstd//:zstd",
     ] + select({
diff --git a/benchmark/run_bench.sh b/benchmark/run_bench.sh
index c35cb898..c014ac1b 100644
--- a/benchmark/run_bench.sh
+++ b/benchmark/run_bench.sh
@@ -1,13 +1,13 @@
 #!/usr/bin/bash
 #
 # 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.
diff --git a/docs/development/index.rst b/docs/development/index.rst
index ebc3047e..fb899cd9 100644
--- a/docs/development/index.rst
+++ b/docs/development/index.rst
@@ -17,3 +17,4 @@ This part contains design principles of SPU and staging PSI documentation.
    psi_protocol_intro
    pir
    add_protocols
+   ir_dump
diff --git a/docs/development/ir_dump.rst b/docs/development/ir_dump.rst
new file mode 100644
index 00000000..194cd2ea
--- /dev/null
+++ b/docs/development/ir_dump.rst
@@ -0,0 +1,174 @@
+Dump IR to DAG
+========================
+
+.. warning::
+    The configuration for dump IR to DAG maybe unstable. Please refer to the :spu_code_host:`spu.proto </spu/blob/main/libspu/spu.proto>` for latest configurations.
+
+Introduction
+------------
+This document provides the demo for how to dump the IR (Intermediate Representation, generated by `XLA <https://www.tensorflow.org/xla/architecture>`_) to a DAG. 
+With the aid of visualized DAG, the execution logic and required operators will be more explicit.
+
+
+.. Note::
+    Please refer to :ref:`/development/compiler.rst` for description of the role of XLA in SPU.
+
+
+TL;DR
+------------
+In case you just want to have a try on obtaining the DAG for executed code, we provide a :spu_code_host:`demo </spu/blob/main/examples/python/ir_dump/ir_dump.py>` that demonstrates the required code modifications to enable dumping the IR of executed code to a custom path.
+
+For those users who have designated requirements (e.g., dump to txt, dot, html), we recommend to read the following detailed step-by-step elaborations.
+
+
+Configuration
+-------------
+Please first have a look at the :spu_code_host:`spu.proto </spu/blob/main/libspu/spu.proto>` for SPU. To dump the IR, we should modify the compiler options, which is shown in the following code snippet.
+
+.. code-block:: protobuf
+    :caption: Compiler Options
+
+    message CompilerOptions {
+        // Pretty print
+        bool enable_pretty_print = 1;
+        string pretty_print_dump_dir = 2;
+        XLAPrettyPrintKind xla_pp_kind = 3;
+        
+        // Disable sqrt(x) + eps to sqrt(x+eps) rewrite
+        bool disable_sqrt_plus_epsilon_rewrite = 10;
+
+        // Disable x/sqrt(y) to x*rsqrt(y) rewrite
+        bool disable_div_sqrt_rewrite = 11;
+
+        // Disable reduce truncation optimization
+        bool disable_reduce_truncation_optimization = 12;
+
+        // Disable maxpooling optimization
+        bool disable_maxpooling_optimization = 13;
+
+        // Disallow mix type operations
+        bool disallow_mix_types_opts = 14;
+
+        // Disable SelectOp optimization
+        bool disable_select_optimization = 15;
+
+        // Enable optimize x/bcast(y) -> x * bcast(1/y)
+        bool enable_optimize_denominator_with_broadcast = 16;
+    }
+
+The part of the above code related to dumping IR to DAG is as follows.
+
+.. code-block:: protobuf
+    :caption: Configurations related to dump IR
+
+    // Pretty print
+    bool enable_pretty_print = 1;
+    string pretty_print_dump_dir = 2;
+    XLAPrettyPrintKind xla_pp_kind = 3;  
+
+In general, we require to focus on three variables to control the dump behavior.
+
+- **enable_pretty_print** is a *bool* value, which denotes whether to dump the IR or not.
+- **pretty_print_dump_dir** is a *string* value, which denotes the dump path (should be a directory).
+- **xla_pp_kind** is a *int* value, of which the range is [0, 1, 2], with each one representing one dump format. To date, we support three kinds of formats: TEXT, DOT and HTML. If you want to obtain the DAG, you should use DOT or HTML.
+
+.. Note::
+    For **DOT** files, you should use `GraphViz <https://graphviz.org/>`_ to convert them to PDF or PNG to visualize the DAG. 
+    
+    While for **HTML** files, you can directly open the them in your Web Browser, which shall render the DAG.
+
+.. code-block:: protobuf
+    :caption: XLA Pretty Print Kind
+
+    enum XLAPrettyPrintKind {
+        TEXT = 0;
+        DOT = 1;
+        HTML = 2;
+    }
+
+
+Pass custom compiler options
+----------------------------------------
+We hereby describe how to manually pass the custom compiler options to dump the IR of executed code.
+
+First of all, we declare an CompilerOptions object. Note that the **pretty_print_dump_dir** is better to be an absolute path.
+
+.. code-block:: python
+    :caption: Declare CompilerOptions object
+
+    import spu.spu_pb2 as spu_pb2
+    copts = spu_pb2.CompilerOptions()
+    copts.enable_pretty_print = True
+    copts.pretty_print_dump_dir = "your/custom/path"
+    copts.xla_pp_kind = 2
+
+Then we pass the CompilerOptions to the executed SPU code.
+
+
+The code shall be modified from 
+
+.. code-block:: python
+    :caption: SPU execution without customized compiler options
+
+    res_spu = ppd.device("SPU")(func)(x_spu, y_spu)
+
+to
+
+.. code-block:: python
+    :caption: SPU execution with customized compiler options
+
+    res_spu = ppd.device("SPU")(func, copts=copts)(x_spu, y_spu)
+
+.. Note::
+    Here, `func` is a Python function. Please refer to the :spu_code_host:`demo </spu/blob/main/examples/python/ir_dump/ir_dump.py>` for the context.
+
+In the end, you can just run the target code and the output (e.g., DOT) can be found in **your custom path**.
+
+Example
+-------
+We here provide the code snippet for dumping IR to HTML files. The DAG for the executed function is illustrated in the end.
+
+.. code-block:: python
+    :caption: Code snippet for dumping the IR of func
+
+    def func(x, y):
+        """
+        Any custom function that consists of operators that SPU currently supports.
+        Here, we define a `max` function use jax.numpy.
+        """
+        return jnp.maximum(x, y)
+    
+    def get_data(seed=123):
+        """
+        Any IO function that loads the data.
+        """
+        np.random.seed(seed)
+        data = np.random.randn(3, 4)
+        return data
+    
+    x = get_data(1)
+    y = get_data(2)
+
+    # SPU secure version
+    x_spu = ppd.device("P1")(lambda x: x)(x)
+    y_spu = ppd.device("P2")(lambda x: x)(y)
+
+    import spu.spu_pb2 as spu_pb2
+    copts = spu_pb2.CompilerOptions()
+    copts.enable_pretty_print = True
+    copts.pretty_print_dump_dir = "dump/path"
+    copts.xla_pp_kind = 2 # HTML format
+
+    res_spu = ppd.device("SPU")(func, copts=copts)(x_spu, y_spu)
+
+.. Note::
+    You may find multiple files in the output directory since XLA has mutliple compile passes and generates multiple IRs, with each corresponding to one DAG.
+
+
+The **HTML** output is rendered as follows.
+
+.. figure:: ../imgs/ir_dag.svg
+   :height: 320
+
+   DAG for executed demo function
+
diff --git a/docs/development/pir.rst b/docs/development/pir.rst
index b23b42db..f44a6b99 100644
--- a/docs/development/pir.rst
+++ b/docs/development/pir.rst
@@ -30,6 +30,7 @@ setup phase
 Generate test usage oprf_key.bin
 
 .. code-block:: bash
+
     dd if=/dev/urandom of=oprf_key.bin bs=32 count=1
 
 Start server's terminal.
diff --git a/docs/development/type_system.rst b/docs/development/type_system.rst
index 88b46efa..bf469d7b 100644
--- a/docs/development/type_system.rst
+++ b/docs/development/type_system.rst
@@ -105,7 +105,7 @@ Ops dispatch
 
 As described above, type helps for dispatching, here we use `MUL` instruction as an example.
 
-.. code-block:: python
+.. code-block:: text
 
   %3 = MUL %1 %2
 
@@ -180,7 +180,7 @@ In the above diagram:
 The above idea can be expressed in code like:
 
 .. code-block:: cpp
-  :linenos:
+  :linenos:  
 
   Value i2f(Value); // convert int to fxp
 
@@ -257,7 +257,7 @@ With *type functor*, we have the following op definitions in `mul` dispatch chai
 
 In dispatch phrase, SPU runtime uses type information to select next dispatch op. In this example, `(x:sfxp, y:sfxp)` is applied op `mul`, via pattern matching we got `(V0=SECRET,D0=FXP), (V1=SECRET,D1=FXP)`, and the dispatch stack looks like:
 
-.. code-block:: python
+.. code-block:: text
 
   mul(x:sfxp, y:sfxp)            # dtype dispatch, use D0=FXP, D1=FXP to select fmul
     fmul(x:sfxp, y:sfxp)         # fixed point arithmetic implementation,
diff --git a/docs/imgs/ir_dag.svg b/docs/imgs/ir_dag.svg
new file mode 100644
index 00000000..4a86e10f
--- /dev/null
+++ b/docs/imgs/ir_dag.svg
@@ -0,0 +1,128 @@
+<svg xmlns="http://www.w3.org/2000/svg" xmlns:xlink="http://www.w3.org/1999/xlink" width="100%" height="100%" viewBox="0.00 0.00 588.48 334.64" id="graph">
+<g id="graph0" class="graph" transform="scale(1 1) rotate(0) translate(4 330.6377)">
+<title>G</title>
+<g id="a_graph0"><a xlink:title=" ">
+<polygon fill="#ffffff" stroke="transparent" points="-4,4 -4,-330.6377 584.48,-330.6377 584.48,4 -4,4"></polygon>
+<text text-anchor="start" x="8" y="-311.0377" font-family="Times,serif" font-weight="bold" font-size="14.00" fill="#000000">module_0000.xla_computation_func.0000.HLO passes.after_pipeline-start.before_operand_upcaster</text>
+<text text-anchor="start" x="232.6853" y="-297.0377" font-family="Times,serif" font-weight="bold" font-size="14.00" fill="#000000">Computation main.5</text>
+</a>
+</g>
+<!-- 25863376 -->
+<g id="node1" class="node">
+<title>25863376</title>
+<polygon fill="#ffe0b2" stroke="#cbae82" points="281.66,-290.6377 188.82,-290.6377 188.82,-254.6377 281.66,-254.6377 281.66,-290.6377"></polygon>
+<text text-anchor="start" x="201.6183" y="-275.0377" font-family="Times,serif" font-weight="bold" font-size="14.00" fill="#000000">Parameter 0</text>
+<text text-anchor="start" x="197.0305" y="-261.0377" font-family="Times,serif" font-size="14.00" fill="#000000">f32[3,4]{1,0}</text>
+</g>
+<!-- 25748112 -->
+<g id="node3" class="node">
+<title>25748112</title>
+<g id="a_node3"><a xlink:title="xla_computation(func)/jit(main)/max
+source: /home/haoqi.whq/.cache/bazel/_bazel_haoqi.whq/bd15e67f09906d483a88b2a872e74859/execroot/spulib/bazel-out/k8-opt/bin/examples/python/dump_xla.runfiles/spulib/examples/python/dump_xla.py:59">
+<polygon fill="#ffffff" stroke="#9e9e9e" points="336.66,-218.6377 243.82,-218.6377 243.82,-182.6377 336.66,-182.6377 336.66,-218.6377"></polygon>
+<text text-anchor="start" x="256.6015" y="-203.0377" font-family="Times,serif" font-weight="bold" font-size="14.00" fill="#000000">maximum.3</text>
+<text text-anchor="start" x="252.0305" y="-189.0377" font-family="Times,serif" font-size="14.00" fill="#000000">f32[3,4]{1,0}</text>
+</a>
+</g>
+</g>
+<!-- 25863376&#45;&gt;25748112 -->
+<g id="edge1" class="edge">
+<title>25863376-&gt;25748112</title>
+<g id="a_edge1"><a xlink:title="Arg_0.1 -> maximum.3">
+<path fill="none" stroke="#000000" d="M249.1188,-254.4691C255.4536,-246.1763 263.0784,-236.1947 270.0326,-227.091"></path>
+<polygon fill="none" stroke="#000000" points="272.8852,-229.1223 276.1743,-219.051 267.3225,-224.873 272.8852,-229.1223"></polygon>
+</a>
+</g>
+<text text-anchor="middle" x="271.8879" y="-234.3862" font-family="Times,serif" font-size="14.00" fill="#000000">0</text>
+</g>
+<!-- 26469072 -->
+<g id="node2" class="node">
+<title>26469072</title>
+<polygon fill="#ffe0b2" stroke="#cbae82" points="391.66,-290.6377 298.82,-290.6377 298.82,-254.6377 391.66,-254.6377 391.66,-290.6377"></polygon>
+<text text-anchor="start" x="311.6183" y="-275.0377" font-family="Times,serif" font-weight="bold" font-size="14.00" fill="#000000">Parameter 1</text>
+<text text-anchor="start" x="307.0305" y="-261.0377" font-family="Times,serif" font-size="14.00" fill="#000000">f32[3,4]{1,0}</text>
+</g>
+<!-- 26469072&#45;&gt;25748112 -->
+<g id="edge2" class="edge">
+<title>26469072-&gt;25748112</title>
+<g id="a_edge2"><a xlink:title="Arg_1.2 -> maximum.3">
+<path fill="none" stroke="#000000" d="M331.3612,-254.4691C325.0264,-246.1763 317.4016,-236.1947 310.4474,-227.091"></path>
+<polygon fill="none" stroke="#000000" points="313.1575,-224.873 304.3057,-219.051 307.5948,-229.1223 313.1575,-224.873"></polygon>
+</a>
+</g>
+<text text-anchor="middle" x="322.0258" y="-224.1244" font-family="Times,serif" font-size="14.00" fill="#000000">1</text>
+</g>
+<!-- 18620560 -->
+<g id="node4" class="node">
+<title>18620560</title>
+<polygon fill="#ffffff" stroke="#9e9e9e" points="325.5433,-146.6377 254.9367,-146.6377 254.9367,-110.6377 325.5433,-110.6377 325.5433,-146.6377"></polygon>
+<text text-anchor="start" x="270.9921" y="-131.0377" font-family="Times,serif" font-weight="bold" font-size="14.00" fill="#000000">tuple.4</text>
+<text text-anchor="start" x="262.8385" y="-117.0377" font-family="Times,serif" font-size="14.00" fill="#000000">(f32[3,4])</text>
+</g>
+<!-- 25748112&#45;&gt;18620560 -->
+<g id="edge3" class="edge">
+<title>25748112-&gt;18620560</title>
+<g id="a_edge3"><a xlink:title="maximum.3 -> tuple.4">
+<path fill="none" stroke="#000000" d="M290.24,-182.4691C290.24,-174.7687 290.24,-165.612 290.24,-157.0543"></path>
+<polygon fill="none" stroke="#000000" points="293.7401,-157.0509 290.24,-147.051 286.7401,-157.051 293.7401,-157.0509"></polygon>
+</a>
+</g>
+</g>
+<!-- cluster_24950960 -->
+<g id="node5" class="node">
+<title>cluster_24950960</title>
+<g id="a_node5"><a xlink:title=" ">
+<ellipse fill="#bcaaa4" stroke="#8c7b75" cx="290.24" cy="-37.3188" rx="37.1386" ry="37.1386"></ellipse>
+<text text-anchor="start" x="271.1888" y="-33.1188" font-family="Times,serif" font-size="14.00" fill="#000000">ROOT</text>
+</a>
+</g>
+</g>
+<!-- 18620560&#45;&gt;cluster_24950960 -->
+<g id="edge4" class="edge">
+<title>18620560-&gt;cluster_24950960</title>
+<g id="a_edge4"><a xlink:title=" ">
+<path fill="none" stroke="#000000" d="M290.24,-110.596C290.24,-103.1939 290.24,-94.2407 290.24,-85.1718"></path>
+<polygon fill="#000000" stroke="#000000" points="293.7401,-84.96 290.24,-74.9601 286.7401,-84.9601 293.7401,-84.96"></polygon>
+</a>
+</g>
+</g>
+</g>
+<style>
+  @import url(https://fonts.googleapis.com/css?family=Roboto:400,700);
+  svg text {
+    font-family: 'Roboto';
+    font-size: 12px;
+  }
+
+    #node4:hover ~ #edge4 text { fill: #1976d2; }
+  #node4:hover ~ #edge4 path { stroke: #1976d2; stroke-width: .2em; }
+  #node4:hover ~ #edge4 polygon { fill: #1976d2; stroke: #1976d2; stroke-width: .2em; }
+
+  #node5:hover ~ #edge4 text { fill: #d32f2f; }
+  #node5:hover ~ #edge4 path { stroke: #d32f2f; stroke-width: .2em; }
+  #node5:hover ~ #edge4 polygon { fill: #d32f2f; stroke: #d32f2f; stroke-width: .2em; }
+
+  #node3:hover ~ #edge3 text { fill: #1976d2; }
+  #node3:hover ~ #edge3 path { stroke: #1976d2; stroke-width: .2em; }
+  #node3:hover ~ #edge3 polygon { fill: #1976d2; stroke: #1976d2; stroke-width: .2em; }
+
+  #node4:hover ~ #edge3 text { fill: #d32f2f; }
+  #node4:hover ~ #edge3 path { stroke: #d32f2f; stroke-width: .2em; }
+  #node4:hover ~ #edge3 polygon { fill: #d32f2f; stroke: #d32f2f; stroke-width: .2em; }
+
+  #node2:hover ~ #edge2 text { fill: #1976d2; }
+  #node2:hover ~ #edge2 path { stroke: #1976d2; stroke-width: .2em; }
+  #node2:hover ~ #edge2 polygon { fill: #1976d2; stroke: #1976d2; stroke-width: .2em; }
+
+  #node3:hover ~ #edge2 text { fill: #d32f2f; }
+  #node3:hover ~ #edge2 path { stroke: #d32f2f; stroke-width: .2em; }
+  #node3:hover ~ #edge2 polygon { fill: #d32f2f; stroke: #d32f2f; stroke-width: .2em; }
+
+  #node1:hover ~ #edge1 text { fill: #1976d2; }
+  #node1:hover ~ #edge1 path { stroke: #1976d2; stroke-width: .2em; }
+  #node1:hover ~ #edge1 polygon { fill: #1976d2; stroke: #1976d2; stroke-width: .2em; }
+
+  #node3:hover ~ #edge1 text { fill: #d32f2f; }
+  #node3:hover ~ #edge1 path { stroke: #d32f2f; stroke-width: .2em; }
+  #node3:hover ~ #edge1 polygon { fill: #d32f2f; stroke: #d32f2f; stroke-width: .2em; }
+</style></svg>
\ No newline at end of file
diff --git a/docs/reference/np_op_status.json b/docs/reference/np_op_status.json
index 5031f749..830a0ca5 100644
--- a/docs/reference/np_op_status.json
+++ b/docs/reference/np_op_status.json
@@ -1,1 +1,1 @@
-[{"name": "abs", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "add", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "angle", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "angle", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arccos", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arccosh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "nan"}, {"name": "arcsin", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arcsinh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "arctan", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arctan2", "dtypes": ["float32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arctanh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "[-1, 1] nan"}, {"name": "argmax", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "argmin", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "array_equal", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "array_equiv", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "atleast_1d", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "atleast_2d", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "atleast_3d", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "bitwise_and", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "bitwise_not", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "bitwise_or", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "bitwise_xor", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "cbrt", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.cbrt"}, {"name": "ceil", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "conj", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "conjugate", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "copysign", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "shift"}, {"name": "cos", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.cosine"}, {"name": "cosh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "deg2rad", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "divide", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "divmod", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "ediff1d", "dtypes": ["int32"], "status": "Status.Pass", "note": ""}, {"name": "equal", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "exp", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "exp2", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "expm1", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "fabs", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "fix", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "float_power", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "floor", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "floor_divide", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "fmax", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "fmin", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "fmod", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "gcd", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "secret while"}, {"name": "greater", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "greater_equal", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "heaviside", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "hypot", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "i0", "dtypes": ["float32"], "status": "Status.Failed", "note": "accuracy"}, {"name": "imag", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "invert", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "isclose", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "iscomplex", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "isfinite", "dtypes": ["float32"], "status": "Status.UnSupport", "note": ""}, {"name": "isinf", "dtypes": ["float32"], "status": "Status.UnSupport", "note": ""}, {"name": "isnan", "dtypes": ["float32"], "status": "Status.UnSupport", "note": ""}, {"name": "isneginf", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "isposinf", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "isreal", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "isrealobj", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "kron", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "lcm", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "secret while"}, {"name": "ldexp", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Failed", "note": "IEEE-754"}, {"name": "left_shift", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "less", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "less_equal", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log10", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log1p", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log1p", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log2", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "logaddexp", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "logaddexp2", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "logical_and", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "logical_not", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "logical_or", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "logical_xor", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "maximum", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "mean", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "minimum", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "mod", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "modf", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "multiply", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "nanargmax", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.UnSupport", "note": ""}, {"name": "nanargmin", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.UnSupport", "note": ""}, {"name": "nanmean", "dtypes": ["float32"], "status": "Status.UnSupport", "note": ""}, {"name": "nanprod", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.UnSupport", "note": ""}, {"name": "nansum", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.UnSupport", "note": ""}, {"name": "negative", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "nextafter", "dtypes": ["float32"], "status": "Status.Failed", "note": "IEEE-754"}, {"name": "not_equal", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "outer", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "polyval", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "positive", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "power", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "prod", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "rad2deg", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "ravel", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "real", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "reciprocal", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "remainder", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "right_shift", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "rint", "dtypes": ["float32"], "status": "Status.SysError", "note": "stablehlo.round_nearest_even"}, {"name": "sign", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "signbit", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "sin", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.sine"}, {"name": "sinc", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.sine"}, {"name": "sinh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.cosine"}, {"name": "sqrt", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "square", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "subtract", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "sum", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "tan", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.sine"}, {"name": "tanh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "transpose", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "true_divide", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "trunc", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "unwrap", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}]
\ No newline at end of file
+[{"name": "abs", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "add", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "angle", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "angle", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arccos", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arccosh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "nan"}, {"name": "arcsin", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arcsinh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "arctan", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arctan2", "dtypes": ["float32"], "status": "Status.SysError", "note": "stablehlo.atan2"}, {"name": "arctanh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "[-1, 1] nan"}, {"name": "argmax", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "argmin", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "array_equal", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "array_equiv", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "atleast_1d", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "atleast_2d", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "atleast_3d", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "bitwise_and", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "bitwise_not", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "bitwise_or", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "bitwise_xor", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "cbrt", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.cbrt"}, {"name": "ceil", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "conj", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "conjugate", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "copysign", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "shift"}, {"name": "cos", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.cosine"}, {"name": "cosh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "deg2rad", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "divide", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "divmod", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "ediff1d", "dtypes": ["int32"], "status": "Status.Pass", "note": ""}, {"name": "equal", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "exp", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "exp2", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "expm1", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "fabs", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "fix", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "float_power", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "floor", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "floor_divide", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "fmax", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "fmin", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "fmod", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "gcd", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "secret while"}, {"name": "greater", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "greater_equal", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "heaviside", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "hypot", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "i0", "dtypes": ["float32"], "status": "Status.Failed", "note": "accuracy"}, {"name": "imag", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "invert", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "isclose", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "iscomplex", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "isfinite", "dtypes": ["float32"], "status": "Status.UnSupport", "note": ""}, {"name": "isinf", "dtypes": ["float32"], "status": "Status.UnSupport", "note": ""}, {"name": "isnan", "dtypes": ["float32"], "status": "Status.UnSupport", "note": ""}, {"name": "isneginf", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "isposinf", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "isreal", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "isrealobj", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "kron", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "lcm", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "secret while"}, {"name": "ldexp", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Failed", "note": "IEEE-754"}, {"name": "left_shift", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "less", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "less_equal", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log10", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log1p", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log1p", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "log2", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "logaddexp", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "logaddexp2", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "logical_and", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "logical_not", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "logical_or", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "logical_xor", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "maximum", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "mean", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "mean", "dtypes": ["float32"], "status": "Status.PassNoGen", "note": ""}, {"name": "minimum", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "mod", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "modf", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "multiply", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "nanargmax", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.UnSupport", "note": ""}, {"name": "nanargmin", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.UnSupport", "note": ""}, {"name": "nanmean", "dtypes": ["float32"], "status": "Status.UnSupport", "note": ""}, {"name": "nanprod", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.UnSupport", "note": ""}, {"name": "nansum", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.UnSupport", "note": ""}, {"name": "negative", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "nextafter", "dtypes": ["float32"], "status": "Status.Failed", "note": "IEEE-754"}, {"name": "not_equal", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "outer", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "polyval", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "positive", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "power", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "prod", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "rad2deg", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "ravel", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "real", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "reciprocal", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}, {"name": "remainder", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "right_shift", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "rint", "dtypes": ["float32"], "status": "Status.SysError", "note": "stablehlo.round_nearest_even"}, {"name": "sign", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "signbit", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "sin", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.sine"}, {"name": "sinc", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.sine"}, {"name": "sinh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.cosine"}, {"name": "sqrt", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "square", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "subtract", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "sum", "dtypes": ["int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "tan", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.SysError", "note": "stablehlo.sine"}, {"name": "tanh", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "transpose", "dtypes": ["bool_", "float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "true_divide", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "trunc", "dtypes": ["float32", "int16", "int32", "uint16", "uint32"], "status": "Status.Pass", "note": ""}, {"name": "unwrap", "dtypes": ["float32"], "status": "Status.Pass", "note": ""}]
\ No newline at end of file
diff --git a/docs/reference/np_op_status.md b/docs/reference/np_op_status.md
index 826480d9..7f5b6757 100644
--- a/docs/reference/np_op_status.md
+++ b/docs/reference/np_op_status.md
@@ -18,11 +18,11 @@ Please contact us if
 # Tested Operators List
 
 ## abs
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.abs.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -34,11 +34,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## add
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.add.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -50,66 +50,66 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## angle
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.angle.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.atan2
 ## angle
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.angle.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.atan2
 ## arccos
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.arccos.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.atan2
 ## arccosh
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.arccosh.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 nan
 ## arcsin
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.arcsin.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.atan2
 ## arcsinh
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.arcsinh.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -120,44 +120,44 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## arctan
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.arctan.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.atan2
 ## arctan2
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.arctan2.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.atan2
 ## arctanh
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.arctanh.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 [-1, 1] nan
 ## argmax
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.argmax.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -169,11 +169,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## argmin
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.argmin.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -185,11 +185,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## array_equal
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.array_equal.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -200,11 +200,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## array_equiv
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.array_equiv.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -215,11 +215,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## atleast_1d
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.atleast_1d.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -230,11 +230,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## atleast_2d
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.atleast_2d.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -245,11 +245,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## atleast_3d
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.atleast_3d.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -260,11 +260,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## bitwise_and
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.bitwise_and.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -274,11 +274,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## bitwise_not
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.bitwise_not.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -288,11 +288,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## bitwise_or
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.bitwise_or.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -302,11 +302,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## bitwise_xor
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.bitwise_xor.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -316,22 +316,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## cbrt
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.cbrt.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.cbrt
 ## ceil
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.ceil.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -342,11 +342,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## conj
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.conj.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -357,11 +357,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## conjugate
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.conjugate.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -372,33 +372,33 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## copysign
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.copysign.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 shift
 ## cos
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.cos.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.cosine
 ## cosh
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.cosh.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -409,22 +409,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## deg2rad
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.deg2rad.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - float32
 
 ## divide
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.divide.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -435,11 +435,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## divmod
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.divmod.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -450,22 +450,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## ediff1d
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.ediff1d.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - int32
 
 ## equal
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.equal.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -477,11 +477,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## exp
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.exp.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -492,11 +492,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## exp2
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.exp2.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -507,11 +507,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## expm1
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.expm1.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -522,22 +522,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## fabs
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.fabs.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - float32
 
 ## fix
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.fix.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -548,22 +548,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## float_power
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.float_power.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - float32
 
 ## floor
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.floor.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -574,11 +574,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## floor_divide
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.floor_divide.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -589,11 +589,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## fmax
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.fmax.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -604,11 +604,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## fmin
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.fmin.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -619,11 +619,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## fmod
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.fmod.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -634,22 +634,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## gcd
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.gcd.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 secret while
 ## greater
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.greater.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -661,11 +661,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## greater_equal
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.greater_equal.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -677,11 +677,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## heaviside
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.heaviside.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -692,11 +692,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## hypot
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.hypot.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -707,22 +707,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## i0
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.i0.html
 ### Status
-  
-Result is incorrect.  
+
+Result is incorrect.
 Please check *Note* for details.
 ### Note
 
 
 accuracy
 ## imag
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.imag.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -733,11 +733,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## invert
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.invert.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -747,11 +747,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## isclose
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.isclose.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -763,11 +763,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## iscomplex
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.iscomplex.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -778,54 +778,54 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## isfinite
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.isfinite.html
 ### Status
-  
-Not supported by design. We couldn't fix this in near future.  
+
+Not supported by design. We couldn't fix this in near future.
 Please check *Note* for details.
 ## isinf
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.isinf.html
 ### Status
-  
-Not supported by design. We couldn't fix this in near future.  
+
+Not supported by design. We couldn't fix this in near future.
 Please check *Note* for details.
 ## isnan
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.isnan.html
 ### Status
-  
-Not supported by design. We couldn't fix this in near future.  
+
+Not supported by design. We couldn't fix this in near future.
 Please check *Note* for details.
 ## isneginf
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.isneginf.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - float32
 
 ## isposinf
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.isposinf.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - float32
 
 ## isreal
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.isreal.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -836,11 +836,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## isrealobj
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.isrealobj.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -851,11 +851,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## kron
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.kron.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -866,33 +866,33 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## lcm
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.lcm.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 secret while
 ## ldexp
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.ldexp.html
 ### Status
-  
-Result is incorrect.  
+
+Result is incorrect.
 Please check *Note* for details.
 ### Note
 
 
 IEEE-754
 ## left_shift
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.left_shift.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -902,11 +902,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## less
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.less.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -918,11 +918,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## less_equal
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.less_equal.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -934,11 +934,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## log
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.log.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -949,11 +949,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## log10
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.log10.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -964,11 +964,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## log1p
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.log1p.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -979,11 +979,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## log1p
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.log1p.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -994,11 +994,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## log2
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.log2.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1009,33 +1009,33 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## logaddexp
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.logaddexp.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - float32
 
 ## logaddexp2
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.logaddexp2.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - float32
 
 ## logical_and
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.logical_and.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1047,11 +1047,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## logical_not
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.logical_not.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1063,11 +1063,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## logical_or
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.logical_or.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1079,11 +1079,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## logical_xor
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.logical_xor.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1095,11 +1095,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## maximum
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.maximum.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1111,11 +1111,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## mean
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.mean.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1126,11 +1126,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## minimum
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.minimum.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1142,11 +1142,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## mod
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.mod.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1157,11 +1157,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## modf
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.modf.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1172,11 +1172,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## multiply
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.multiply.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1188,46 +1188,46 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## nanargmax
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.nanargmax.html
 ### Status
-  
-Not supported by design. We couldn't fix this in near future.  
+
+Not supported by design. We couldn't fix this in near future.
 Please check *Note* for details.
 ## nanargmin
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.nanargmin.html
 ### Status
-  
-Not supported by design. We couldn't fix this in near future.  
+
+Not supported by design. We couldn't fix this in near future.
 Please check *Note* for details.
 ## nanmean
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.nanmean.html
 ### Status
-  
-Not supported by design. We couldn't fix this in near future.  
+
+Not supported by design. We couldn't fix this in near future.
 Please check *Note* for details.
 ## nanprod
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.nanprod.html
 ### Status
-  
-Not supported by design. We couldn't fix this in near future.  
+
+Not supported by design. We couldn't fix this in near future.
 Please check *Note* for details.
 ## nansum
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.nansum.html
 ### Status
-  
-Not supported by design. We couldn't fix this in near future.  
+
+Not supported by design. We couldn't fix this in near future.
 Please check *Note* for details.
 ## negative
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.negative.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1238,22 +1238,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## nextafter
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.nextafter.html
 ### Status
-  
-Result is incorrect.  
+
+Result is incorrect.
 Please check *Note* for details.
 ### Note
 
 
 IEEE-754
 ## not_equal
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.not_equal.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1265,11 +1265,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## outer
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.outer.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1280,11 +1280,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## polyval
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.polyval.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1295,11 +1295,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## positive
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.positive.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1310,11 +1310,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## power
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.power.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1325,11 +1325,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## prod
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.prod.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1341,22 +1341,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## rad2deg
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.rad2deg.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - float32
 
 ## ravel
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.ravel.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1368,11 +1368,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## real
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.real.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1383,22 +1383,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## reciprocal
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.reciprocal.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
 - float32
 
 ## remainder
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.remainder.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1409,11 +1409,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## right_shift
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.right_shift.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1423,22 +1423,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## rint
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.rint.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.round_nearest_even
 ## sign
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.sign.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1449,11 +1449,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## signbit
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.signbit.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1464,44 +1464,44 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## sin
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.sin.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.sine
 ## sinc
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.sinc.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.sine
 ## sinh
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.sinh.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.cosine
 ## sqrt
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.sqrt.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1512,11 +1512,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## square
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.square.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1527,11 +1527,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## subtract
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.subtract.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1542,11 +1542,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## sum
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.sum.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1556,22 +1556,22 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## tan
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.tan.html
 ### Status
-  
-Not supported by compiler or runtime. But we could implement on demand in future.  
+
+Not supported by compiler or runtime. But we could implement on demand in future.
 Please check *Note* for details.
 ### Note
 
 
 stablehlo.sine
 ## tanh
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.tanh.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1582,11 +1582,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## transpose
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.transpose.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1598,11 +1598,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## true_divide
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.true_divide.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1613,11 +1613,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## trunc
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.trunc.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
@@ -1628,11 +1628,11 @@ Please check *Supported Dtypes* as well.
 - uint32
 
 ## unwrap
-  
+
 JAX NumPy Document link: https://jax.readthedocs.io/en/latest/_autosummary/jax.numpy.unwrap.html
 ### Status
-  
-**PASS**  
+
+**PASS**
 Please check *Supported Dtypes* as well.
 ### Supported Dtypes
 
diff --git a/docs/reference/pphlo_op_doc.md b/docs/reference/pphlo_op_doc.md
index ebbb09ec..4b78db6b 100644
--- a/docs/reference/pphlo_op_doc.md
+++ b/docs/reference/pphlo_op_doc.md
@@ -323,7 +323,7 @@ See https://www.tensorflow.org/xla/operation_semantics#concatenate.
 
 Traits: AlwaysSpeculatableImplTrait, SameOperandsAndResultElementType
 
-Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
+Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)
 
 Effects: MemoryEffects::Effect{}
 
@@ -553,7 +553,7 @@ See https://www.tensorflow.org/xla/operation_semantics#dynamicslice.
 
 Traits: AlwaysSpeculatableImplTrait
 
-Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
+Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)
 
 Effects: MemoryEffects::Effect{}
 
@@ -587,7 +587,7 @@ See https://www.tensorflow.org/xla/operation_semantics#dynamicupdateslice.
 
 Traits: AlwaysSpeculatableImplTrait
 
-Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
+Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)
 
 Effects: MemoryEffects::Effect{}
 
@@ -1434,7 +1434,7 @@ Effects: MemoryEffects::Effect{}
 
 
     The `pphlo.return` operation terminates a region and returns values.
-  
+
 
 
 Traits: AlwaysSpeculatableImplTrait, Terminator
diff --git a/docs/reference/psi_config.md b/docs/reference/psi_config.md
index 85e3ca94..929e9d54 100644
--- a/docs/reference/psi_config.md
+++ b/docs/reference/psi_config.md
@@ -11,13 +11,13 @@
     - [MemoryPsiConfig](#memorypsiconfig)
     - [OutputParams](#outputparams)
     - [PsiResultReport](#psiresultreport)
-  
+
 
 
 - Enums
     - [CurveType](#curvetype)
     - [PsiType](#psitype)
-  
+
 
 
 - [Scalar Value Types](#scalar-value-types)
diff --git a/docs/reference/runtime_config.md b/docs/reference/runtime_config.md
index 73584994..bf5f42fa 100644
--- a/docs/reference/runtime_config.md
+++ b/docs/reference/runtime_config.md
@@ -5,14 +5,15 @@
 
 
 - Messages
+    - [CompilationSource](#compilationsource)
+    - [CompilerOptions](#compileroptions)
     - [ExecutableProto](#executableproto)
     - [RuntimeConfig](#runtimeconfig)
     - [ShapeProto](#shapeproto)
     - [TTPBeaverConfig](#ttpbeaverconfig)
     - [ValueMeta](#valuemeta)
     - [ValueProto](#valueproto)
-    - [XlaMeta](#xlameta)
-  
+
 
 
 - Enums
@@ -24,8 +25,10 @@
     - [RuntimeConfig.ExpMode](#runtimeconfigexpmode)
     - [RuntimeConfig.LogMode](#runtimeconfiglogmode)
     - [RuntimeConfig.SigmoidMode](#runtimeconfigsigmoidmode)
+    - [SourceIRType](#sourceirtype)
     - [Visibility](#visibility)
-  
+    - [XLAPrettyPrintKind](#xlaprettyprintkind)
+
 
 
 - [Scalar Value Types](#scalar-value-types)
@@ -37,6 +40,39 @@
 ## Messages
 
 
+### CompilationSource
+
+
+
+| Field | Type | Description |
+| ----- | ---- | ----------- |
+| ir_type | [ SourceIRType](#sourceirtype) | Input IR type |
+| ir_txt | [ bytes](#bytes) | IR |
+| input_visibility | [repeated Visibility](#visibility) | Input visibilities |
+ <!-- end Fields -->
+ <!-- end HasFields -->
+
+
+### CompilerOptions
+
+
+
+| Field | Type | Description |
+| ----- | ---- | ----------- |
+| enable_pretty_print | [ bool](#bool) | Pretty print |
+| pretty_print_dump_dir | [ string](#string) | none |
+| xla_pp_kind | [ XLAPrettyPrintKind](#xlaprettyprintkind) | none |
+| disable_sqrt_plus_epsilon_rewrite | [ bool](#bool) | Disable sqrt(x) + eps to sqrt(x+eps) rewrite |
+| disable_div_sqrt_rewrite | [ bool](#bool) | Disable x/sqrt(y) to x*rsqrt(y) rewrite |
+| disable_reduce_truncation_optimization | [ bool](#bool) | Disable reduce truncation optimization |
+| disable_maxpooling_optimization | [ bool](#bool) | Disable maxpooling optimization |
+| disallow_mix_types_opts | [ bool](#bool) | Disallow mix type operations |
+| disable_select_optimization | [ bool](#bool) | Disable SelectOp optimization |
+| enable_optimize_denominator_with_broadcast | [ bool](#bool) | Enable optimize x/bcast(y) -> x * bcast(1/y) |
+ <!-- end Fields -->
+ <!-- end HasFields -->
+
+
 ### ExecutableProto
 The executable format accepted by SPU runtime.
 
@@ -74,7 +110,7 @@ The SPU runtime configuration.
 | ----- | ---- | ----------- |
 | protocol | [ ProtocolKind](#protocolkind) | The protocol kind. |
 | field | [ FieldType](#fieldtype) | The field type. |
-| fxp_fraction_bits | [ int64](#int64) | Number of fraction bits of fixed-point number. |
+| fxp_fraction_bits | [ int64](#int64) | Number of fraction bits of fixed-point number. 0(default) indicates implementation defined. |
 | enable_action_trace | [ bool](#bool) | When enabled, runtime prints verbose info of the call stack, debug purpose only. |
 | enable_type_checker | [ bool](#bool) | When enabled, runtime checks runtime type infos against the compile-time ones, exceptions are raised if mismatches happen. Note: Runtime outputs prefer runtime type infos even when flag is on. |
 | enable_pphlo_trace | [ bool](#bool) | When enabled, runtime prints executed pphlo list, debug purpose only. |
@@ -150,17 +186,6 @@ The spu Value proto, used for spu value serialization.
 | content | [ bytes](#bytes) | The runtime/protocol dependent value data. |
  <!-- end Fields -->
  <!-- end HasFields -->
-
-
-### XlaMeta
-Internal representation used by compiler
-
-
-| Field | Type | Description |
-| ----- | ---- | ----------- |
-| inputs | [repeated Visibility](#visibility) | none |
- <!-- end Fields -->
- <!-- end HasFields -->
  <!-- end messages -->
 
 ## Enums
@@ -290,6 +315,18 @@ The sigmoid approximation method.
 
 
 
+### SourceIRType
+Compiler relate definition
+////////////////////////////////////////////////////////////////////////
+
+| Name | Number | Description |
+| ---- | ------ | ----------- |
+| XLA | 0 | none |
+| MLIR_HLO | 1 | none |
+
+
+
+
 ### Visibility
 The visibility type.
 
@@ -304,6 +341,18 @@ performance significantly.
 | VIS_PUBLIC | 2 | Visible(public) for all parties. |
 
 
+
+
+### XLAPrettyPrintKind
+
+
+| Name | Number | Description |
+| ---- | ------ | ----------- |
+| TEXT | 0 | none |
+| DOT | 1 | none |
+| HTML | 2 | none |
+
+
  <!-- end Enums -->
  <!-- end Files -->
 
diff --git a/docs/reference/update_complexity_doc.sh b/docs/reference/update_complexity_doc.sh
index 3777c51a..bdd256d7 100755
--- a/docs/reference/update_complexity_doc.sh
+++ b/docs/reference/update_complexity_doc.sh
@@ -20,6 +20,6 @@ SCRIPTPATH=`dirname $SCRIPT`
 REPOPATH=`realpath $SCRIPTPATH/../../`
 
 bazel run //libspu/mpc/tools:complexity -- --out="$REPOPATH/docs/reference/complexity.json"
-python $REPOPATH/docs/reference/gen_complexity_md.py --in="$REPOPATH/docs/reference/complexity.json" --out="$REPOPATH/docs/reference/complexity.md" 
+python $REPOPATH/docs/reference/gen_complexity_md.py --in="$REPOPATH/docs/reference/complexity.json" --out="$REPOPATH/docs/reference/complexity.md"
 
 # echo $REPOPATH
diff --git a/docs/reference/update_np_op_status_doc.sh b/docs/reference/update_np_op_status_doc.sh
index 11a5863b..a70b98d2 100644
--- a/docs/reference/update_np_op_status_doc.sh
+++ b/docs/reference/update_np_op_status_doc.sh
@@ -20,5 +20,5 @@ SCRIPTPATH=`dirname $SCRIPT`
 REPOPATH=`realpath $SCRIPTPATH/../../`
 
 bazel run //spu/tests:np_op_status -- --out="$REPOPATH/docs/reference/np_op_status.json"
-python $REPOPATH/docs/reference/gen_np_op_status_doc.py --in="$REPOPATH/docs/reference/np_op_status.json" --out="$REPOPATH/docs/reference/np_op_status.md" 
+python $REPOPATH/docs/reference/gen_np_op_status_doc.py --in="$REPOPATH/docs/reference/np_op_status.json" --out="$REPOPATH/docs/reference/np_op_status.md"
 
diff --git a/docs/reference/xla_status.md b/docs/reference/xla_status.md
index 067f22fb..dd4844aa 100644
--- a/docs/reference/xla_status.md
+++ b/docs/reference/xla_status.md
@@ -145,11 +145,11 @@ Count: Total = 2, fully supported = 2
 | `clamp`        | fully                       |
 | `concatenate`  | fully                       |
 | `collective_permute` | no                    |
-| `convolution`  | fully                       | 
+| `convolution`  | fully                       |
 | `copy`         | no                          |
 | `cross-replica-sum` | no                     |
 | `custom_call`  | no                          |
-| `dot`          | fully                       | 
+| `dot`          | fully                       |
 | `dot_general`  | fully                       |
 | `einsum`       | fully                       |
 | `unary_einsum` | fully                       |
diff --git a/examples/cpp/pir/keyword_pir_client.cc b/examples/cpp/pir/keyword_pir_client.cc
index 064ce6bc..2ed7dd3c 100644
--- a/examples/cpp/pir/keyword_pir_client.cc
+++ b/examples/cpp/pir/keyword_pir_client.cc
@@ -17,7 +17,7 @@
 // > bazel build //examples/cpp/pir:keyword_pir_client -c opt
 //
 // To run the example, start terminals:
-// > ./keyword_pir_client -rank 1 -in_path ../../data/psi_client_data.csv.csv 
+// > ./keyword_pir_client -rank 1 -in_path ../../data/psi_client_data.csv.csv
 // >       -key_columns id -out_path pir_out.csv
 // clang-format on
 
diff --git a/examples/cpp/pir/keyword_pir_server.cc b/examples/cpp/pir/keyword_pir_server.cc
index b3b2817c..48ecdcb9 100644
--- a/examples/cpp/pir/keyword_pir_server.cc
+++ b/examples/cpp/pir/keyword_pir_server.cc
@@ -17,7 +17,7 @@
 // > bazel build //examples/cpp/pir:keyword_pir_server -c opt
 //
 // To run the example, start terminals:
-// > ./keyword_pir_server -rank 0 -setup_path pir_setup_dir  
+// > ./keyword_pir_server -rank 0 -setup_path pir_setup_dir
 // >        -oprf_key_path secret_key.bin
 // clang-format on
 
diff --git a/examples/cpp/pir/keyword_pir_setup.cc b/examples/cpp/pir/keyword_pir_setup.cc
index 321b3a93..055f2aee 100644
--- a/examples/cpp/pir/keyword_pir_setup.cc
+++ b/examples/cpp/pir/keyword_pir_setup.cc
@@ -20,8 +20,8 @@
 // > dd if=/dev/urandom of=secret_key.bin bs=32 count=1
 //
 // To run the example, start terminals:
-// > ./keyword_pir_setup -in_path ../../data/psi_server_data.csv -oprf_key_path secret_key.bin  
-// >     -key_columns id -label_columns label -data_per_query 1 -label_max_len 40  
+// > ./keyword_pir_setup -in_path ../../data/psi_server_data.csv -oprf_key_path secret_key.bin
+// >     -key_columns id -label_columns label -data_per_query 1 -label_max_len 40
 // >     -setup_path pir_setup_dir
 // clang-format on
 
diff --git a/examples/cpp/simple_dp_psi.cc b/examples/cpp/simple_dp_psi.cc
index 13bb0e3c..108cf17c 100644
--- a/examples/cpp/simple_dp_psi.cc
+++ b/examples/cpp/simple_dp_psi.cc
@@ -17,7 +17,7 @@
 // > bazel build //examples/cpp:simple_dp_psi -c opt
 //
 // To run the example, start two terminals:
-// > ./simple_dp_psi -rank 0 -in_path examples/data/psi_1.csv -field_names id 
+// > ./simple_dp_psi -rank 0 -in_path examples/data/psi_1.csv -field_names id
 // > ./simple_dp_psi -rank 1 -in_path examples/data/psi_2.csv -field_names id -out_path /tmp/p2.out
 // clang-format on
 
diff --git a/examples/cpp/simple_lr.cc b/examples/cpp/simple_lr.cc
index 85179127..67320f93 100644
--- a/examples/cpp/simple_lr.cc
+++ b/examples/cpp/simple_lr.cc
@@ -32,13 +32,14 @@
 #include "libspu/device/io.h"
 #include "libspu/kernel/hal/hal.h"
 #include "libspu/kernel/hal/type_cast.h"
+#include "libspu/mpc/factory.h"
 
 using namespace spu::kernel;
 
 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});
+  auto padding = hal::constant(ctx, 1.0F, spu::DT_F32, {x.shape()[0], 1});
   auto padded_x = hal::concatenate(ctx, {x, hal::seal(ctx, padding)}, 1);
   auto pred = hal::logistic(ctx, hal::matmul(ctx, padded_x, w));
 
@@ -49,9 +50,9 @@ spu::Value train_step(spu::SPUContext* ctx, const spu::Value& x,
   auto grad = hal::matmul(ctx, hal::transpose(ctx, padded_x), err);
 
   SPDLOG_DEBUG("[SSLR] Step = LR / B * Grad");
-  auto lr = hal::constant(ctx, 0.0001F, spu::DT_FXP);
+  auto lr = hal::constant(ctx, 0.0001F, spu::DT_F32);
   auto msize =
-      hal::constant(ctx, static_cast<float>(y.shape()[0]), spu::DT_FXP);
+      hal::constant(ctx, static_cast<float>(y.shape()[0]), spu::DT_F32);
   auto p1 = hal::mul(ctx, lr, hal::reciprocal(ctx, msize));
   auto step = hal::mul(ctx, hal::broadcast_to(ctx, p1, grad.shape()), grad);
 
@@ -64,7 +65,7 @@ spu::Value train_step(spu::SPUContext* ctx, const spu::Value& x,
 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});
+  auto w = hal::constant(ctx, 0.0F, spu::DT_F32, {x.shape()[1] + 1, 1});
 
   // Run train loop
   for (size_t epoch = 0; epoch < num_epoch; ++epoch) {
@@ -89,7 +90,7 @@ spu::Value train(spu::SPUContext* ctx, const spu::Value& x, const spu::Value& y,
 
 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 padding = hal::constant(ctx, 1.0F, spu::DT_F32, {x.shape()[0], 1});
   auto padded_x = hal::concatenate(ctx, {x, hal::seal(ctx, padding)}, 1);
   return hal::matmul(ctx, padded_x, weight);
 }
@@ -168,6 +169,8 @@ int main(int argc, char** argv) {
 
   auto sctx = MakeSPUContext();
 
+  spu::mpc::Factory::RegisterProtocol(sctx.get(), sctx->lctx());
+
   const auto& [x, y] = infeed(sctx.get(), ds, HasLabel.getValue());
 
   const auto w =
diff --git a/examples/cpp/simple_psi.cc b/examples/cpp/simple_psi.cc
index 1ab8f687..a9787fa3 100644
--- a/examples/cpp/simple_psi.cc
+++ b/examples/cpp/simple_psi.cc
@@ -14,7 +14,7 @@
 
 // clang-format off
 // To run the example, start two terminals:
-// > bazel run //examples/cpp:simple_psi -- -rank 0 -protocol 1 -in_path examples/data/psi_1.csv -field_names id -out_path /tmp/p1.out 
+// > bazel run //examples/cpp:simple_psi -- -rank 0 -protocol 1 -in_path examples/data/psi_1.csv -field_names id -out_path /tmp/p1.out
 // > bazel run //examples/cpp:simple_psi -- -rank 1 -protocol 1 -in_path examples/data/psi_2.csv -field_names id -out_path /tmp/p2.out
 // clang-format on
 
diff --git a/examples/cpp/utils.cc b/examples/cpp/utils.cc
index 829301f0..96809b0c 100644
--- a/examples/cpp/utils.cc
+++ b/examples/cpp/utils.cc
@@ -17,10 +17,12 @@
 #include "absl/strings/match.h"
 #include "absl/strings/str_split.h"
 
+#include "libspu/core/config.h"
+
 #include "libspu/spu.pb.h"
 
 llvm::cl::opt<std::string> Parties(
-    "parties", llvm::cl::init("127.0.0.1:9530,127.0.0.1:9531"),
+    "parties", llvm::cl::init("127.0.0.1:39530,127.0.0.1:39531"),
     llvm::cl::desc("server list, format: host1:port1[,host2:port2, ...]"));
 llvm::cl::opt<uint32_t> Rank("rank", llvm::cl::init(0),
                              llvm::cl::desc("self rank"));
@@ -53,7 +55,10 @@ std::unique_ptr<spu::SPUContext> MakeSPUContext() {
   config.set_protocol(static_cast<spu::ProtocolKind>(ProtocolKind.getValue()));
   config.set_field(static_cast<spu::FieldType>(Field.getValue()));
 
+  populateRuntimecConfig(config);
+
   config.set_enable_action_trace(EngineTrace.getValue());
   config.set_enable_type_checker(EngineTrace.getValue());
+
   return std::make_unique<spu::SPUContext>(config, lctx);
 }
diff --git a/examples/python/ir_dump/BUILD.bazel b/examples/python/ir_dump/BUILD.bazel
new file mode 100644
index 00000000..15736655
--- /dev/null
+++ b/examples/python/ir_dump/BUILD.bazel
@@ -0,0 +1,28 @@
+# 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.
+
+load("@rules_python//python:defs.bzl", "py_binary")
+
+package(default_visibility = ["//visibility:public"])
+
+py_binary(
+    name = "ir_dump",
+    srcs = ["ir_dump.py"],
+    data = [
+        "//examples/python/conf",
+    ],
+    deps = [
+        "//spu/utils:distributed",
+    ],
+)
diff --git a/examples/python/ir_dump/README.md b/examples/python/ir_dump/README.md
new file mode 100644
index 00000000..a6664cb0
--- /dev/null
+++ b/examples/python/ir_dump/README.md
@@ -0,0 +1,18 @@
+# How to dump IR to DAG in SPU
+
+In this directory, we provide an example of how to dump the IR (Intermediate Representation, generated by [XLA](https://www.tensorflow.org/xla/architecture))
+of executed code to a DAG (Directed Acyclic Graph) that illistrates the computation graph.
+
+Please refer to [spu.proto](../../../libspu/spu.proto) for more descriptions of the compiler options for dumping IR.
+
+1. Launch SPU backend runtime
+
+    ``` shell
+    bazel run -c opt //examples/python/utils:nodectl -- up
+    ```
+
+2. Run `ir_dump` example
+
+    ``` shell
+    bazel run -c opt //examples/python/ir_dump:ir_dump
+    ```
diff --git a/examples/python/ir_dump/ir_dump.py b/examples/python/ir_dump/ir_dump.py
new file mode 100644
index 00000000..fd081657
--- /dev/null
+++ b/examples/python/ir_dump/ir_dump.py
@@ -0,0 +1,93 @@
+# 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.
+
+
+# Start nodes.
+# > bazel run -c opt //examples/python/utils:nodectl -- up
+#
+# Run this example script.
+# > bazel run -c opt //examples/python/ir_dump:ir_dump
+
+
+import argparse
+import json
+import logging
+import os
+
+import jax.numpy as jnp
+import numpy as np
+
+import spu.utils.distributed as ppd
+import spu.spu_pb2 as spu_pb2
+
+logging.basicConfig(level=logging.INFO)
+
+parser = argparse.ArgumentParser(description="distributed driver.")
+parser.add_argument("-c", "--config", default="examples/python/conf/3pc.json")
+parser.add_argument("-d", "--dir", default="ppdump")
+args = parser.parse_args()
+
+with open(args.config, "r") as file:
+    conf = json.load(file)
+
+ppd.init(conf["nodes"], conf["devices"])
+
+
+# enable dump ir to custom path
+dump_path = os.path.join(os.path.expanduser("~"), args.dir)
+logging.info(f"Dump path: {dump_path}")
+# refer to spu.proto for more detailed configuration
+copts = spu_pb2.CompilerOptions()
+copts.enable_pretty_print = True
+copts.pretty_print_dump_dir = dump_path
+copts.xla_pp_kind = 2
+
+
+def func(x, y):
+    """
+    Any custom function that consists of operators that SPU currently supports.
+    Here, we define a `max` function use jax.numpy.
+    """
+    return jnp.maximum(x, y)
+
+
+def get_data(seed=123):
+    """
+    Any IO function that loads the data.
+    """
+    np.random.seed(seed)
+    data = np.random.randn(3, 4)
+    return data
+
+
+def main():
+    # CPU plaintext version
+    x = get_data(1)
+    y = get_data(2)
+
+    res = func(x, y)
+    logging.info(f"\nx: {x}\ny: {y}\nres: {res}")
+
+    # SPU secure version
+    x_spu = ppd.device("P1")(lambda x: x)(x)
+    y_spu = ppd.device("P2")(lambda x: x)(y)
+
+    res_spu = ppd.device("SPU")(func, copts=copts)(x_spu, y_spu)
+    logging.info(f"\nx: {ppd.get(x_spu)}\ny: {ppd.get(y_spu)}\nres: {ppd.get(res_spu)}")
+
+    assert jnp.allclose(res, ppd.get(res_spu), 1e-4)
+
+
+if __name__ == "__main__":
+    main()
diff --git a/examples/python/ml/README.md b/examples/python/ml/README.md
index f32f5b5e..09cc861c 100644
--- a/examples/python/ml/README.md
+++ b/examples/python/ml/README.md
@@ -4,17 +4,23 @@ This directory contains examples demonstrating how to use SPU to write privacy-p
 
 * [ss_lr](ss_lr/): Private training of linear/logistic regression models with JAX.
 * [ss_xgb](ss_xgb/): Private training of an XGB model with JAX.
-* [jax_lr](jax_lr/): Private training of a logistic regression model with JAX (different from `ss_lr`, we run a complete JAX training function in this demo).
+* [jax_lr](jax_lr/): Private training of a logistic regression model with JAX (different from `ss_lr`,
+                     we run a complete JAX training function in this demo).
 * [jax_svm](jax_svm/): Private training of an SVM model with JAX.
 * [jax_kmeans](jax_kmeans/): Private training of K-Means clustering with JAX.
 * [flax_mlp](flax_mlp/): Private training of an MLP model with [Flax](https://github.com/google/flax).
-* [stax_nn](stax_nn/): Private training of four widely-evaluated neural networks for MNIST classification with [Stax](https://jax.readthedocs.io/en/latest/jax.example_libraries.stax.html).
-* [stax_mnist_classifier](stax_mnist_classifier/): Private training of a simple neural network for MNIST classification with [Stax](https://jax.readthedocs.io/en/latest/jax.example_libraries.stax.html).
+* [stax_nn](stax_nn/): Private training of four widely-evaluated neural networks for MNIST classification with
+                       [Stax](https://jax.readthedocs.io/en/latest/jax.example_libraries.stax.html).
+* [stax_mnist_classifier](stax_mnist_classifier/): Private training of a simple neural network for MNIST classification with
+                                                   [Stax](https://jax.readthedocs.io/en/latest/jax.example_libraries.stax.html).
 * [flax_resnet](flax_resnet/): Private training of a [ResNet](https://arxiv.org/abs/1512.03385) model with [Flax](https://github.com/google/flax) library.
-* [flax_gpt2](flax_gpt2/): Private inference of a pre-trained [GPT2](https://cdn.openai.com/better-language-models/language_models_are_unsupervised_multitask_learners.pdf) model with [Flax](https://github.com/google/flax) library.
-* [flax_vae](flax_vae/): Private training of a [VAE](http://arxiv.org/abs/1312.6114) model for MNIST digit reconstruction with [Flax](https://github.com/google/flax) library.
+* [flax_gpt2](flax_gpt2/): Private inference of a pre-trained
+                           [GPT2](https://cdn.openai.com/better-language-models/language_models_are_unsupervised_multitask_learners.pdf)
+                           model with [Flax](https://github.com/google/flax) library.
+* [flax_vae](flax_vae/): Private training of a [VAE](http://arxiv.org/abs/1312.6114) model for MNIST digit reconstruction with
+                         [Flax](https://github.com/google/flax) library.
 * [haiku_lstm](haiku_lstm/): Private training of an LSTM model with [Haiku](https://github.com/deepmind/dm-haiku).
-* [jraph_gnn](jraph_gnn/): Private training of a [graph convolutional network](https://arxiv.org/abs/1609.02907) model with [Jraph](https://github.com/deepmind/jraph).
+* [jraph_gnn](jraph_gnn/): Private training of a [graph convolutional network](https://arxiv.org/abs/1609.02907) model with
+                           [Jraph](https://github.com/deepmind/jraph).
 * [tf_experiment](tf_experiment/): Private training of a logistic regression model with TensorFlow (**experimental**).
 * [torch_experiment](torch_experiment/): Private inference of a linear regression model with PyTorch (**experimental**).
-
diff --git a/examples/python/ml/flax_gpt2/README.md b/examples/python/ml/flax_gpt2/README.md
index 724ecd4f..7eb25fc4 100644
--- a/examples/python/ml/flax_gpt2/README.md
+++ b/examples/python/ml/flax_gpt2/README.md
@@ -1,16 +1,22 @@
-This example demonstrates how to use SPU to run private inference on a pre-trained [GPT2](https://cdn.openai.com/better-language-models/language_models_are_unsupervised_multitask_learners.pdf) model.
+# Flax GPT2 Example
+
+This example demonstrates how to use SPU to run private inference on a pre-trained
+[GPT2](https://cdn.openai.com/better-language-models/language_models_are_unsupervised_multitask_learners.pdf) model.
 
 1. Install huggingface transformers library
-    ```
+
+    ```sh
     pip install 'transformers[flax]'
     ```
 
 2. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- --config `pwd`/examples/python/ml/flax_gpt2/3pc.json up
     ```
 
 3. Run `flax_gpt2` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/flax_gpt2 -- --config `pwd`/examples/python/ml/flax_gpt2/3pc.json
     ```
diff --git a/examples/python/ml/flax_mlp/README.md b/examples/python/ml/flax_mlp/README.md
index ad7a0f97..74607495 100644
--- a/examples/python/ml/flax_mlp/README.md
+++ b/examples/python/ml/flax_mlp/README.md
@@ -1,11 +1,15 @@
+# Flax MLP Example
+
 This example demonstrates how to use SPU to train an MLP model privately.
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `flax_mlp` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/flax_mlp
     ```
diff --git a/examples/python/ml/flax_resnet/README.md b/examples/python/ml/flax_resnet/README.md
index 5564c3f6..a834188c 100644
--- a/examples/python/ml/flax_resnet/README.md
+++ b/examples/python/ml/flax_resnet/README.md
@@ -1,15 +1,19 @@
+# Flax ResNet Example
+
 This example demonstrates how to use SPU to train the [ResNet](https://arxiv.org/abs/1512.03385) model privately.
 
 This example comes from Flax official github repo:
 
-https://github.com/google/flax/tree/main/examples/imagenet
+<https://github.com/google/flax/tree/main/examples/imagenet>
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- --config `pwd`/examples/python/ml/flax_resnet/3pc.json up
     ```
 
 2. Run `flax_resnet` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/flax_resnet -- --config `pwd`/examples/python/ml/flax_resnet/3pc.json --num_epochs 5
     ```
diff --git a/examples/python/ml/flax_vae/README.md b/examples/python/ml/flax_vae/README.md
index 24172ea6..f6ac599e 100644
--- a/examples/python/ml/flax_vae/README.md
+++ b/examples/python/ml/flax_vae/README.md
@@ -1,16 +1,20 @@
+# Flax VAE Example
+
 This example demonstrates how to use SPU to train the [VAE](http://arxiv.org/abs/1312.6114) model privately.
 
 This example comes from Flax official github repo:
 
-https://github.com/google/flax/tree/main/examples/vae
+<https://github.com/google/flax/tree/main/examples/vae>
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `flax_vae` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/flax_vae -- --output_dir `pwd` --num_epochs 5
     ```
 
diff --git a/examples/python/ml/haiku_lstm/README.md b/examples/python/ml/haiku_lstm/README.md
index 282a5ff6..4a150d4b 100644
--- a/examples/python/ml/haiku_lstm/README.md
+++ b/examples/python/ml/haiku_lstm/README.md
@@ -1,21 +1,26 @@
+# haiku LSTM Example
+
 This example demonstrates how to use SPU to train an LSTM model privately.
 
 This example comes from Haiku official github repo:
 
-https://github.com/deepmind/dm-haiku/blob/main/examples/haiku_lstms.ipynb
+<https://github.com/deepmind/dm-haiku/blob/main/examples/haiku_lstms.ipynb>
 
 1. Install dependencies
-    ```
+
+    ```sh
     pip install -r requirements.txt
     ```
 
 2. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 3. Run `haiku_lstm` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/haiku_lstm -- --output_dir `pwd`
     ```
 
diff --git a/examples/python/ml/jax_kmeans/README.md b/examples/python/ml/jax_kmeans/README.md
index 0c3fa11b..a59cde85 100644
--- a/examples/python/ml/jax_kmeans/README.md
+++ b/examples/python/ml/jax_kmeans/README.md
@@ -1,11 +1,15 @@
+# Jax Kmeans Example
+
 This example demonstrates how to use SPU to train K-Means clustering privately.
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `jax_kmeans` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/jax_kmeans
     ```
diff --git a/examples/python/ml/jax_lr/README.md b/examples/python/ml/jax_lr/README.md
index 4842d05b..fc0e6580 100644
--- a/examples/python/ml/jax_lr/README.md
+++ b/examples/python/ml/jax_lr/README.md
@@ -1,11 +1,15 @@
+# Jax LR Example
+
 This example demonstrates how to use SPU to train a logistic regression model privately.
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `jax_lr` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/jax_lr
     ```
diff --git a/examples/python/ml/jax_svm/README.md b/examples/python/ml/jax_svm/README.md
index 0bca4891..175ac07b 100644
--- a/examples/python/ml/jax_svm/README.md
+++ b/examples/python/ml/jax_svm/README.md
@@ -1,11 +1,15 @@
+# Jax SVM Example
+
 This example demonstrates how to use SPU to train an SVM model privately.
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `jax_svm` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/jax_svm
     ```
diff --git a/examples/python/ml/jraph_gnn/README.md b/examples/python/ml/jraph_gnn/README.md
index 65240efd..0739ddd4 100644
--- a/examples/python/ml/jraph_gnn/README.md
+++ b/examples/python/ml/jraph_gnn/README.md
@@ -1,11 +1,14 @@
+# Jraph GNN Example
+
 This example demonstrates how to use SPU to train a Graph Convolutional Network model privately.
 
 This example comes from Jraph official github repo:
 
-https://github.com/deepmind/jraph/blob/master/jraph/examples/zacharys_karate_club.py
+<https://github.com/deepmind/jraph/blob/master/jraph/examples/zacharys_karate_club.py>
 
 1. Install dependencies
-    ```
+
+    ```sh
     pip install -r requirements.txt
     ```
 
@@ -14,15 +17,17 @@ https://github.com/deepmind/jraph/blob/master/jraph/examples/zacharys_karate_clu
     This example requires a higher precision setting than the default.
 
     Set `"fxp_fraction_bits": 24` in `SPU runtime_config`.
-    
+
     The default configuration file locates at [examples/python/conf/3pc.json](../../conf/3pc.json).
 
 3. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 4. Run `jraph_gnn` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/jraph_gnn
     ```
diff --git a/examples/python/ml/ss_lr/README.md b/examples/python/ml/ss_lr/README.md
index bb2e75af..f37d9674 100644
--- a/examples/python/ml/ss_lr/README.md
+++ b/examples/python/ml/ss_lr/README.md
@@ -1,11 +1,15 @@
+# SS-LR Example
+
 This example demonstrates how to use SPU to train linear/logistic regression models privately.
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `ss_lr` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/ss_lr
     ```
diff --git a/examples/python/ml/ss_xgb/README.md b/examples/python/ml/ss_xgb/README.md
index 08ad8c18..ba51376e 100644
--- a/examples/python/ml/ss_xgb/README.md
+++ b/examples/python/ml/ss_xgb/README.md
@@ -1,11 +1,15 @@
+# SS-XGB Example
+
 This example demonstrates how to use SPU to train XGB models privately.
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `ss_xgb` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/ss_xgb
     ```
diff --git a/examples/python/ml/stax_mnist_classifier/README.md b/examples/python/ml/stax_mnist_classifier/README.md
index d3bdb2d9..71556cb7 100644
--- a/examples/python/ml/stax_mnist_classifier/README.md
+++ b/examples/python/ml/stax_mnist_classifier/README.md
@@ -1,11 +1,16 @@
-This example demonstrates how to use SPU to train a neural network model privately for MNIST classification with [Stax](https://jax.readthedocs.io/en/latest/jax.example_libraries.stax.html) library.
+# Stax MNIST Classifier Example
+
+This example demonstrates how to use SPU to train a neural network model privately for MNIST classification with
+[Stax](https://jax.readthedocs.io/en/latest/jax.example_libraries.stax.html) library.
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `stax_mnist_classifier` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/stax_mnist_classifier
     ```
diff --git a/examples/python/ml/stax_nn/README.md b/examples/python/ml/stax_nn/README.md
index ca2e686b..9c2b190a 100644
--- a/examples/python/ml/stax_nn/README.md
+++ b/examples/python/ml/stax_nn/README.md
@@ -1,13 +1,17 @@
-This example demonstrates how to use SPU to train neural network models privately for image classification. 
+# Stax NN Example
+
+This example demonstrates how to use SPU to train neural network models privately for image classification.
 
 These models are widely used for evaluation benchmarks in MPC-enabled literature such as [deep-mpc](https://arxiv.org/abs/2107.00501) and [SeureNN](https://eprint.iacr.org/2018/442.pdf).
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `stax_nn` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/stax_nn
     ```
diff --git a/examples/python/ml/tf_experiment/README.md b/examples/python/ml/tf_experiment/README.md
index df356055..c56d9690 100644
--- a/examples/python/ml/tf_experiment/README.md
+++ b/examples/python/ml/tf_experiment/README.md
@@ -1,13 +1,17 @@
-This example demonstrates how to use SPU to train a logistic regression privately with TensorFlow. 
+# TensorFlow Example
+
+This example demonstrates how to use SPU to train a logistic regression privately with TensorFlow.
 
 Currently, SPU's support of TensorFlow is **experimental**.
 
 1. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 2. Run `tf_experiment` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/tf_experiment
     ```
diff --git a/examples/python/ml/torch_experiment/README.md b/examples/python/ml/torch_experiment/README.md
index 67a0d903..dbc728c2 100644
--- a/examples/python/ml/torch_experiment/README.md
+++ b/examples/python/ml/torch_experiment/README.md
@@ -1,19 +1,24 @@
-This example demonstrates how to use SPU to make inferences on a linear regression model privately with PyTorch. 
+# Torch Example
+
+This example demonstrates how to use SPU to make inferences on a linear regression model privately with PyTorch.
 
 The model is trained with plaintext publicly. Currently, SPU's support of PyTorch is **experimental** and we only tested on Linux.
 
 1. Install a third-party dependency [Torch-MLIR](https://github.com/llvm/torch-mlir).
-    ```
+
+    ```sh
     pip install https://github.com/llvm/torch-mlir/releases/download/snapshot-20220830.581/torch-1.13.0.dev20220830+cpu-cp38-cp38-linux_x86_64.whl
     pip install https://github.com/llvm/torch-mlir/releases/download/snapshot-20220830.581/torch_mlir-20220830.581-cp38-cp38-linux_x86_64.whl
     ```
 
 2. Launch SPU backend runtime
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/utils:nodectl -- up
     ```
 
 3. Run `torch_experiment` example
-    ```
+
+    ```sh
     bazel run -c opt //examples/python/ml/torch_experiment
     ```
diff --git a/libspu/compiler/common/compilation_context.h b/libspu/compiler/common/compilation_context.h
index de039d21..25e1b59e 100644
--- a/libspu/compiler/common/compilation_context.h
+++ b/libspu/compiler/common/compilation_context.h
@@ -47,6 +47,10 @@ class CompilationContext {
 
   bool hasPrettyPrintEnabled() const { return options_.enable_pretty_print(); }
 
+  XLAPrettyPrintKind getXlaPrettyPrintKind() const {
+    return options_.xla_pp_kind();
+  }
+
   std::filesystem::path getPrettyPrintDir() const;
 
 private:
diff --git a/libspu/compiler/core/core.cc b/libspu/compiler/core/core.cc
index 7651c09d..a9680a7c 100644
--- a/libspu/compiler/core/core.cc
+++ b/libspu/compiler/core/core.cc
@@ -61,6 +61,10 @@ void Core::buildPipeline(mlir::PassManager *pm) {
 
   optPM.addPass(mlir::pphlo::createExpandSecretGatherPass());
 
+  if (options.enable_optimize_denominator_with_broadcast()) {
+    optPM.addPass(mlir::pphlo::createOptimizeDenominatorWithBroadcast());
+  }
+
   optPM.addPass(mlir::createCSEPass());
 
   if (!options.disable_reduce_truncation_optimization()) {
diff --git a/libspu/compiler/front_end/hlo_importer.cc b/libspu/compiler/front_end/hlo_importer.cc
index 5c5c8a47..d22a1923 100644
--- a/libspu/compiler/front_end/hlo_importer.cc
+++ b/libspu/compiler/front_end/hlo_importer.cc
@@ -150,7 +150,20 @@ HloImporter::parseXlaModuleFromString(const std::string &content) {
   if (context_->hasPrettyPrintEnabled()) {
     debug_options.set_xla_dump_hlo_pass_re(".*");
     debug_options.set_xla_dump_to(context_->getPrettyPrintDir().string());
-    debug_options.set_xla_dump_hlo_as_text(true);
+    switch (context_->getXlaPrettyPrintKind()) {
+    case spu::XLAPrettyPrintKind::DOT: {
+      debug_options.set_xla_dump_hlo_as_dot(true);
+      break;
+    }
+    case spu::XLAPrettyPrintKind::HTML: {
+      debug_options.set_xla_dump_hlo_as_html(true);
+      break;
+    }
+    default: {
+      debug_options.set_xla_dump_hlo_as_text(true);
+      break;
+    }
+    }
     debug_options.set_xla_detailed_logging_and_dumping(true);
   }
 
diff --git a/libspu/compiler/passes/BUILD.bazel b/libspu/compiler/passes/BUILD.bazel
index a88347d3..c04e7159 100644
--- a/libspu/compiler/passes/BUILD.bazel
+++ b/libspu/compiler/passes/BUILD.bazel
@@ -222,6 +222,18 @@ spu_cc_library(
     ],
 )
 
+spu_cc_library(
+    name = "optimize_denominator_with_broadcast",
+    srcs = ["optimize_denominator_with_broadcast.cc"],
+    hdrs = ["passes.h"],
+    deps = [
+        ":pass_details",
+        "//libspu/dialect:pphlo_dialect",
+        "@llvm-project//mlir:IR",
+        "@llvm-project//mlir:TransformUtils",
+    ],
+)
+
 spu_cc_library(
     name = "all_passes",
     hdrs = ["register_passes.h"],
@@ -232,6 +244,7 @@ spu_cc_library(
         ":hlo_legalize_to_pphlo",
         ":lower_conversion_cast",
         ":lower_mixed_type_op",
+        ":optimize_denominator_with_broadcast",
         ":optimize_maxpool",
         ":optimize_select",
         ":optimize_sqrt_plus_eps",
diff --git a/libspu/compiler/passes/optimize_denominator_with_broadcast.cc b/libspu/compiler/passes/optimize_denominator_with_broadcast.cc
new file mode 100644
index 00000000..1a63a5b1
--- /dev/null
+++ b/libspu/compiler/passes/optimize_denominator_with_broadcast.cc
@@ -0,0 +1,74 @@
+// 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 "mlir/IR/PatternMatch.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
+
+#include "libspu/compiler/passes/pass_details.h"
+#include "libspu/dialect/pphlo_ops.h"
+
+namespace mlir::pphlo {
+
+namespace {
+
+struct DivRewriter : public OpRewritePattern<DivOp> {
+  explicit DivRewriter(MLIRContext *context)
+      : OpRewritePattern<DivOp>(context) {}
+
+  LogicalResult matchAndRewrite(DivOp op,
+                                PatternRewriter &rewriter) const override {
+    auto denominator = op.getRhs();
+
+    if (auto bcst = denominator.getDefiningOp<BroadcastOp>()) {
+      // denominator is bcast
+      // x/bcast(y) -> x*bcast(1/y)
+      auto bcasted_v = bcst.getOperand();
+      // 1/y
+      auto reci_v = rewriter.create<ReciprocalOp>(
+          op->getLoc(), bcasted_v.getType(), bcasted_v);
+      // bcast(1/y)
+      auto bcsted_reci_v = rewriter.create<BroadcastOp>(
+          op->getLoc(), bcst.getType(), reci_v, bcst.getBroadcastDimensions());
+      // x*bcast(1/y)
+      rewriter.replaceOpWithNewOp<MulOp>(op, op.getType(), op.getLhs(),
+                                         bcsted_reci_v);
+      return success();
+    }
+    return failure();
+  }
+};
+
+struct OptimizeDenominatorWithBcast
+    : public OptimizeDenominatorWithBcastBase<OptimizeDenominatorWithBcast> {
+  void runOnOperation() override {
+    RewritePatternSet patterns(&getContext());
+    populateOwningPatterns(&patterns, &getContext());
+    (void)applyPatternsAndFoldGreedily(getOperation(), std::move(patterns));
+  }
+
+private:
+  static void populateOwningPatterns(RewritePatternSet *patterns,
+                                     MLIRContext *ctx) {
+    patterns->insert<DivRewriter>(ctx);
+  }
+};
+} // namespace
+
+std::unique_ptr<OperationPass<func::FuncOp>>
+createOptimizeDenominatorWithBroadcast() {
+  return std::make_unique<OptimizeDenominatorWithBcast>();
+}
+
+} // namespace mlir::pphlo
diff --git a/libspu/compiler/passes/passes.h b/libspu/compiler/passes/passes.h
index 40be508e..c8eeeea7 100644
--- a/libspu/compiler/passes/passes.h
+++ b/libspu/compiler/passes/passes.h
@@ -67,6 +67,10 @@ std::unique_ptr<OperationPass<func::FuncOp>> createExpandSecretGatherPass();
 // Rewrite x/sqrt(x+eps) -> x*rsqrt(x+eps)
 std::unique_ptr<OperationPass<func::FuncOp>> createRewriteDivSqrtPatterns();
 
+// Optimize x/broadcast(y) into x*broadcast(1/y)
+std::unique_ptr<OperationPass<func::FuncOp>>
+createOptimizeDenominatorWithBroadcast();
+
 } // namespace pphlo
 
 } // namespace mlir
diff --git a/libspu/compiler/passes/passes.td b/libspu/compiler/passes/passes.td
index 9b9353a3..4292b5c5 100644
--- a/libspu/compiler/passes/passes.td
+++ b/libspu/compiler/passes/passes.td
@@ -86,3 +86,9 @@ def ExpandSecretGather: Pass<"expand-secret-gather", "func::FuncOp"> {
   let constructor = "createExpandSecretGatherPass()";
   let dependentDialects = ["pphlo::PPHloDialect"];
 }
+
+def OptimizeDenominatorWithBcast: Pass<"optimize-denominator-with-broadcast", "func::FuncOp"> {
+  let summary = "Optimize x/broadcast(y) into x*broadcast(1/y)";
+  let constructor = "createOptimizeDenominatorWithBroadcast()";
+  let dependentDialects = ["pphlo::PPHloDialect"];
+}
diff --git a/libspu/compiler/readme.md b/libspu/compiler/readme.md
index 6b93dfb2..f17aa5b1 100644
--- a/libspu/compiler/readme.md
+++ b/libspu/compiler/readme.md
@@ -1,5 +1,7 @@
 # MLIR-PPHLO: A MLIR based Secure HLO compiler
 
-This implements a self-contained compiler for a privacy preserving linear algebra set of operations inspired by XLA HLO IR using MLIR components. It is designed to provide an compiler engine for SPU.
+This implements a self-contained compiler for a privacy preserving linear algebra set of operations inspired by XLA HLO IR using MLIR components.
+It is designed to provide an compiler engine for SPU.
 
-Coding practice and conventions in this repository follow the [MLIR Developer Guide](https://mlir.llvm.org/getting_started/DeveloperGuide/) in this repo as part of the intent to act as an incubator for technology to upstream.
+Coding practice and conventions in this repository follow the [MLIR Developer Guide](https://mlir.llvm.org/getting_started/DeveloperGuide/) in
+this repo as part of the intent to act as an incubator for technology to upstream.
diff --git a/libspu/core/context.h b/libspu/core/context.h
index b887bcf5..b34853c7 100644
--- a/libspu/core/context.h
+++ b/libspu/core/context.h
@@ -15,6 +15,7 @@
 #pragma once
 
 #include <memory>
+#include <variant>
 
 #include "yacl/link/link.h"
 
@@ -87,14 +88,12 @@ 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
+      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_;
diff --git a/libspu/core/encoding.cc b/libspu/core/encoding.cc
index 08eabf19..441af9ab 100644
--- a/libspu/core/encoding.cc
+++ b/libspu/core/encoding.cc
@@ -161,7 +161,7 @@ ArrayRef decodeFromRing(const ArrayRef& src, DataType in_dtype, size_t fxp_bits,
         pforeach(0, numel, [&](int64_t idx) {
           dst_ptr[idx * dst_stride] = !((src_ptr[idx * src_stride] & 0x1) == 0);
         });
-      } else if (in_dtype == DT_FXP) {
+      } else if (in_dtype == DT_F32 || in_dtype == DT_F64) {
         const T kScale = T(1) << fxp_bits;
         pforeach(0, numel, [&](int64_t idx) {
           dst_ptr[idx * dst_stride] =
diff --git a/libspu/core/encoding.h b/libspu/core/encoding.h
index bc4dc8af..263a40dc 100644
--- a/libspu/core/encoding.h
+++ b/libspu/core/encoding.h
@@ -63,8 +63,8 @@ namespace spu {
   FN(PT_I64, DT_I64)            \
   FN(PT_U64, DT_U64)            \
   FN(PT_BOOL, DT_I1)            \
-  FN(PT_F32, DT_FXP)            \
-  FN(PT_F64, DT_FXP)
+  FN(PT_F32, DT_F32)            \
+  FN(PT_F64, DT_F64)
 
 #define MAP_DTYPE_TO_PTTYPE(FN) \
   FN(DT_I8, PT_I8)              \
@@ -76,7 +76,8 @@ namespace spu {
   FN(DT_I64, PT_I64)            \
   FN(DT_U64, PT_U64)            \
   FN(DT_I1, PT_BOOL)            \
-  FN(DT_FXP, PT_F32)
+  FN(DT_F32, PT_F32)            \
+  FN(DT_F64, PT_F64)
 
 DataType getEncodeType(PtType pt_type);
 
diff --git a/libspu/core/object.h b/libspu/core/object.h
index d462b951..e987b8a1 100644
--- a/libspu/core/object.h
+++ b/libspu/core/object.h
@@ -16,11 +16,8 @@
 
 #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"
 
diff --git a/libspu/core/type_util.h b/libspu/core/type_util.h
index c9da4ba8..94f68804 100644
--- a/libspu/core/type_util.h
+++ b/libspu/core/type_util.h
@@ -51,7 +51,9 @@ std::ostream& operator<<(std::ostream& os, const Visibility& vtype);
   FN(DT_I64, I64, 64)          \
   FN(DT_U64, U64, 64)
 
-#define FOREACH_FXP_DTYPES(FN) FN(DT_FXP, FXP, 64)
+#define FOREACH_FXP_DTYPES(FN) \
+  FN(DT_F32, F32, 32)          \
+  FN(DT_F64, F64, 64)
 
 #define FOREACH_DTYPES(FN) \
   FOREACH_INT_DTYPES(FN)   \
diff --git a/libspu/core/value.h b/libspu/core/value.h
index 410d322f..547ba92e 100644
--- a/libspu/core/value.h
+++ b/libspu/core/value.h
@@ -62,7 +62,6 @@ class Value final {
   // By default, we can only set dtype when it's not set yet(=DT_INVALID),
   // unless we forcely do it.
   Value& setDtype(DataType new_dtype, bool force = false);
-  Value& asFxp() { return setDtype(DT_FXP); }
 
   // Serialize to protobuf.
   ValueProto toProto() const;
diff --git a/libspu/device/readme.md b/libspu/device/README.md
similarity index 69%
rename from libspu/device/readme.md
rename to libspu/device/README.md
index b17878fa..0f31b312 100644
--- a/libspu/device/readme.md
+++ b/libspu/device/README.md
@@ -2,4 +2,5 @@
 
 This implements an interpreter for MLIR-PPHLO.
 
-Coding practice and conventions in this repository follow the [MLIR Developer Guide](https://mlir.llvm.org/getting_started/DeveloperGuide/) in this repo as part of the intent to act as an incubator for technology to upstream.
+Coding practice and conventions in this repository follow the [MLIR Developer Guide](https://mlir.llvm.org/getting_started/DeveloperGuide/)
+in this repo as part of the intent to act as an incubator for technology to upstream.
diff --git a/libspu/device/pphlo/BUILD.bazel b/libspu/device/pphlo/BUILD.bazel
index ae6b6557..dceca77c 100644
--- a/libspu/device/pphlo/BUILD.bazel
+++ b/libspu/device/pphlo/BUILD.bazel
@@ -28,9 +28,11 @@ spu_cc_library(
     ],
 )
 
-spu_cc_test(
-    name = "pphlo_executor_test",
-    srcs = ["pphlo_executor_test.cc"],
+spu_cc_library(
+    name = "pphlo_executor_test_runner",
+    testonly = True,
+    srcs = ["pphlo_executor_test_runner.cc"],
+    hdrs = ["pphlo_executor_test_runner.h"],
     deps = [
         ":pphlo_executor",
         "//libspu/compiler:compile",
@@ -41,6 +43,14 @@ spu_cc_test(
     ],
 )
 
+spu_cc_test(
+    name = "pphlo_executor_test",
+    srcs = ["pphlo_executor_test.cc"],
+    deps = [
+        ":pphlo_executor_test_runner",
+    ],
+)
+
 spu_cc_library(
     name = "pphlo_verifier",
     srcs = ["pphlo_verifier.cc"],
diff --git a/libspu/device/pphlo/pphlo_executor.cc b/libspu/device/pphlo/pphlo_executor.cc
index a0558f1e..b442da46 100644
--- a/libspu/device/pphlo/pphlo_executor.cc
+++ b/libspu/device/pphlo/pphlo_executor.cc
@@ -110,7 +110,14 @@ spu::DataType getDtypeFromMlirType(mlir::Type mlir_ty) {
         SPU_THROW("unsupported int type {}", mlirObjectToString(mlir_ty));
     }
   } else if (auto flp_ty = express_type.dyn_cast<mlir::FloatType>()) {
-    return spu::DT_FXP;
+    switch (flp_ty.getWidth()) {
+      case 32:
+        return spu::DT_F32;
+      case 64:
+        return spu::DT_F64;
+      default:
+        SPU_THROW("unsupported fp type {}", mlirObjectToString(flp_ty));
+    }
   }
   SPU_THROW("invalid type {}", mlirObjectToString(mlir_ty));
 }
@@ -272,7 +279,7 @@ void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
                                  .getSplatValue<mlir::APFloat>()
                                  .convertToFloat());
       auto eps = kernel::hal::dump_public_as<float>(
-          sctx, kernel::hlo::Epsilon(sctx))[0];
+          sctx, kernel::hlo::Epsilon(sctx, DT_F32))[0];
 
       // Amplify eps to 1/(2^(fxp_bits-2))
       // TODO: Maybe make it configurable?
@@ -1081,7 +1088,7 @@ void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
 
 void execute(OpExecutor *executor, SPUContext *sctx, SymbolScope *sscope,
              mlir::pphlo::EpsilonOp &op, const ExecutionOptions &opts) {
-  auto e = kernel::hlo::Epsilon(sctx);
+  auto e = kernel::hlo::Epsilon(sctx, getDtypeFromMlirType(op.getType()));
   auto shape =
       op->getResultTypes()[0].dyn_cast<mlir::RankedTensorType>().getShape();
   addValue(sscope, op.getResult(), kernel::hlo::Broadcast(sctx, e, shape, {}),
diff --git a/libspu/device/pphlo/pphlo_verifier.cc b/libspu/device/pphlo/pphlo_verifier.cc
index a0af24bb..360dd590 100644
--- a/libspu/device/pphlo/pphlo_verifier.cc
+++ b/libspu/device/pphlo/pphlo_verifier.cc
@@ -49,8 +49,10 @@ mlir::Type getElementType(mlir::MLIRContext *mlir_ctx, const spu::Value &v) {
       return mlir::IntegerType::get(mlir_ctx, 64);
     case DT_U64:
       return mlir::IntegerType::get(mlir_ctx, 64, mlir::IntegerType::Unsigned);
-    case DT_FXP:
+    case DT_F32:
       return mlir::Float32Type::get(mlir_ctx);
+    case DT_F64:
+      return mlir::Float64Type::get(mlir_ctx);
     default:
       SPU_THROW("Should not hit");
   }
@@ -265,6 +267,8 @@ UNARY_VERIFIER(RsqrtOp, evalRsqrtOp)
 UNARY_VERIFIER(SqrtOp, evalSqrtOp)
 UNARY_VERIFIER(RoundOp, evalRoundOp)
 UNARY_VERIFIER(SignOp, evalSignOp)
+UNARY_VERIFIER(Log1pOp, evalLog1pOp)
+UNARY_VERIFIER(Expm1Op, evalExpm1Op)
 
 #undef UNARY_VERIFIER
 
@@ -514,18 +518,6 @@ void PPHloVerifier::verify(mlir::pphlo::ReciprocalOp op,
   SPDLOG_WARN("Missing stablehlo interpreter support");
 }
 
-void PPHloVerifier::verify(mlir::pphlo::Log1pOp op,
-                           absl::Span<const spu::Value> operands,
-                           absl::Span<const spu::Value> expected) {
-  SPDLOG_WARN("Missing stablehlo interpreter support");
-}
-
-void PPHloVerifier::verify(mlir::pphlo::Expm1Op op,
-                           absl::Span<const spu::Value> operands,
-                           absl::Span<const spu::Value> expected) {
-  SPDLOG_WARN("Missing stablehlo interpreter support");
-}
-
 void PPHloVerifier::verify(mlir::pphlo::DotOp op,
                            absl::Span<const spu::Value> operands,
                            absl::Span<const spu::Value> expected) {
diff --git a/libspu/dialect/pphlo_attrs.td b/libspu/dialect/pphlo_attrs.td
index f43b2af2..7e919913 100644
--- a/libspu/dialect/pphlo_attrs.td
+++ b/libspu/dialect/pphlo_attrs.td
@@ -25,10 +25,10 @@ def PPHloDim : ArrayRefParameter<"int64_t", "Dimension">;
 def GatherDimensionNumbers : AttrDef<PPHlo_Dialect, "GatherDimensionNumbers"> {
   let mnemonic = "gather";
   let summary = "Attribute that models the dimension information for gather";
-  let parameters = (ins 
-              PPHloDim: $offsetDims, 
-              PPHloDim: $collapsedSliceDims, 
-              PPHloDim: $startIndexMap, 
+  let parameters = (ins
+              PPHloDim: $offsetDims,
+              PPHloDim: $collapsedSliceDims,
+              PPHloDim: $startIndexMap,
               "int64_t": $indexVectorDim);
   let hasCustomAssemblyFormat = 1;
 }
diff --git a/libspu/dialect/pphlo_dialect.td b/libspu/dialect/pphlo_dialect.td
index 26b5daff..236de329 100644
--- a/libspu/dialect/pphlo_dialect.td
+++ b/libspu/dialect/pphlo_dialect.td
@@ -33,7 +33,7 @@ def PPHlo_Dialect : Dialect {
   string description = [{
     PPHLO represents a high level abstraction for language use by SPU.
     It implements a subset of mlir mhlo ops with it's own privacy-preserving focused type system.
-    
+
     Learn more about mlir hlo at https://github.com/tensorflow/mlir-hlo
   }];
   let name = "pphlo";
diff --git a/libspu/dialect/pphlo_ops.cc b/libspu/dialect/pphlo_ops.cc
index 8e95ab18..e605096c 100644
--- a/libspu/dialect/pphlo_ops.cc
+++ b/libspu/dialect/pphlo_ops.cc
@@ -988,6 +988,62 @@ LogicalResult IotaOp::verify() {
   return success();
 }
 
+LogicalResult SliceOp::verify() {
+  auto rankedTy = getOperand().getType();
+  // slice_i2
+  ShapedType attrTy = getStartIndices().getType();
+  if (attrTy.getRank() != 1) {
+    return emitOpError(
+        llvm::formatv("start_indices has rank {0} instead of required rank 1",
+                      attrTy.getRank()));
+  }
+
+  // slice_c2
+  int64_t rank = rankedTy.getRank();
+  if (attrTy.getNumElements() != rank) {
+    return emitOpError(
+        llvm::formatv("the number of elements in start_indices ({0}) does not "
+                      "match the rank of the operand ({1})",
+                      attrTy.getNumElements(), rank));
+  }
+
+  auto start = getStartIndices().getValues<int64_t>();
+  auto limit = getLimitIndices().getValues<int64_t>();
+  auto strideVals = getStrides().getValues<int64_t>();
+
+  for (int64_t i = 0, e = rank; i != e; i++) {
+    // slice_c3
+    if (start[i] < 0) {
+      return emitOpError(llvm::formatv(
+          "negative start index {0} in dimension {1}", start[i], i));
+    }
+
+    int64_t operandSize = rankedTy.getDimSize(i);
+    // slice_c3
+    if (limit[i] > operandSize) {
+      return emitOpError(llvm::formatv(
+          "limit index {0} is larger than dimension size {1} in dimension {2}",
+          limit[i], operandSize, i));
+    }
+
+    // slice_c3
+    if (start[i] > limit[i]) {
+      return emitOpError(llvm::formatv(
+          "start index {0}  is larger than limit index {1} in dimension {2}",
+          start[i], limit[i], i));
+    }
+
+    // slice_c4
+    if (strideVals[i] <= 0) {
+      return emitOpError(
+          llvm::formatv("stride must be positive but got {0} in dimension {1}",
+                        strideVals[i], i));
+    }
+  }
+
+  return success();
+}
+
 LogicalResult inferDynamicSliceOp(std::optional<Location> location,
                                   Type operandType, TypeRange startIndicesTypes,
                                   DenseIntElementsAttr sliceSizes,
@@ -1684,8 +1740,8 @@ ParseResult parseWindowAttributes(OpAsmParser& parser,
       return failure();
     }
 
-    // parse the attribute value. We need to support either 1D and Nx2 array of
-    // integers to parse.
+    // parse the attribute value. We need to support either 1D and Nx2 array
+    // of integers to parse.
     llvm::SmallVector<int64_t> values;
     auto int64_parser = [&]() {
       return parser.parseInteger(values.emplace_back(0));
diff --git a/libspu/dialect/pphlo_ops.td b/libspu/dialect/pphlo_ops.td
index dee9d9c7..6c5e309c 100644
--- a/libspu/dialect/pphlo_ops.td
+++ b/libspu/dialect/pphlo_ops.td
@@ -339,7 +339,7 @@ def PPHLO_MinOp
   let summary = "Minimum operator";
   let description = [{
     Returns `min(lhs, rhs)` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
     }];
@@ -349,7 +349,7 @@ def PPHLO_DivOp : PPHLO_BinaryElementwiseOp<"divide", [Pure]> {
   let summary = "Division operator";
   let description = [{
     Returns `lhs / rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -360,7 +360,7 @@ def PPHLO_AddOp
   let summary = "Addition operator";
   let description = [{
     Returns `lhs + rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -371,7 +371,7 @@ def PPHLO_MulOp
   let summary = "Multiplication operator";
   let description = [{
     Returns `lhs * rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -381,7 +381,7 @@ def PPHLO_PowOp : PPHLO_BinaryElementwiseOp<"power", [Pure]> {
   let summary = "Power operator";
   let description = [{
     Returns `lhs ^ rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -403,7 +403,7 @@ def PPHLO_ShiftLeftOp
   let summary = "Shift Left operator";
   let description = [{
     Returns `lhs << rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -415,7 +415,7 @@ def PPHLO_ShiftRightLogicalOp
   let summary = "Shift right logical operator";
   let description = [{
     Returns logical `lhs >> rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -427,7 +427,7 @@ def PPHLO_ShiftRightArithmeticOp
   let summary = "Shift right arithmetic operator";
   let description = [{
     Returns arithmetic `lhs >> rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -446,7 +446,7 @@ def PPHLO_AndOp : PPHLO_BinaryLogicalElementwiseOp<"and"> {
   let summary = "Logical and";
   let description = [{
     Returns `logical_and(lhs, rhs)` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -456,7 +456,7 @@ def PPHLO_OrOp : PPHLO_BinaryLogicalElementwiseOp<"or"> {
   let summary = "Logical or";
   let description = [{
     Returns `logical_or(lhs, rhs)` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -466,7 +466,7 @@ def PPHLO_XorOp : PPHLO_BinaryLogicalElementwiseOp<"xor"> {
   let summary = "Logical xor";
   let description = [{
     Returns `logical_xor(lhs, rhs)` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations.
   }];
@@ -481,7 +481,7 @@ def PPHLO_EqualOp
   let summary = "Equal comparison operator";
   let description = [{
     Returns `lhs` == `rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_comparison_operations.
   }];
@@ -495,7 +495,7 @@ def PPHLO_NotEqualOp
   let summary = "Not-equal comparison operator";
   let description = [{
     Returns `lhs` != `rhs` element-wise.
-  
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_comparison_operations.
   }];
@@ -509,7 +509,7 @@ def PPHLO_GreaterEqualOp
   let summary = "greater_equal comparison operator";
   let description = [{
     Returns `lhs` >= `rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_comparison_operations.
   }];
@@ -541,7 +541,7 @@ def PPHLO_LessOp
   let summary = "less comparison operator";
   let description = [{
     Returns `lhs` < `rhs` element-wise.
-    
+
     See
     https://www.tensorflow.org/xla/operation_semantics#element-wise_comparison_operations.
   }];
@@ -590,7 +590,7 @@ def PPHLO_BroadcastOp
 
     See https://www.tensorflow.org/xla/broadcasting.
   }];
-  let arguments = (ins PPHLO_Tensor: $operand, 
+  let arguments = (ins PPHLO_Tensor: $operand,
                       I64ElementsAttr: $broadcast_dimensions);
 
   let results = (outs PPHLO_Tensor);
@@ -670,9 +670,9 @@ def PPHLO_SelectOp
 
     See https://www.tensorflow.org/xla/operation_semantics#select.
   }];
-  let arguments = (ins 
-                  PPHLO_IntTensor: $pred, 
-                  PPHLO_Tensor: $on_true, 
+  let arguments = (ins
+                  PPHLO_IntTensor: $pred,
+                  PPHLO_Tensor: $on_true,
                   PPHLO_Tensor: $on_false);
 
   let results = (outs PPHLO_Tensor);
@@ -852,6 +852,8 @@ def PPHLO_SliceOp : PPHLO_Op<"slice", [
                    : $strides);
 
   let results = (outs PPHLO_Tensor);
+
+  let hasVerifier = 1;
 }
 
 def PPHLO_DbgPrintOp : PPHLO_Op<"dbg_print", []> {
diff --git a/libspu/dialect/pphlo_types.td b/libspu/dialect/pphlo_types.td
index 013a56b5..247ad392 100644
--- a/libspu/dialect/pphlo_types.td
+++ b/libspu/dialect/pphlo_types.td
@@ -53,16 +53,16 @@ def PPHLO_SecretType : PPHLO_BaseType<"Secret", "sec"> {
     let summary = "PPHlo secret type";
 }
 
-def PPHLO_PublicIntType : Type<And<[PPHLO_PublicType.predicate, 
+def PPHLO_PublicIntType : Type<And<[PPHLO_PublicType.predicate,
     CPred<"$_self.cast<::mlir::pphlo::PublicType>().getBase().isa<::mlir::IntegerType>()">]>, "public integer type", "::pphlo::PubIntType">;
 
-def PPHLO_SecretIntType : Type<And<[PPHLO_SecretType.predicate, 
+def PPHLO_SecretIntType : Type<And<[PPHLO_SecretType.predicate,
     CPred<"$_self.cast<::mlir::pphlo::SecretType>().getBase().isa<::mlir::IntegerType>()">]>, "secret integer type", "::pphlo::SecIntType">;
 
-def PPHLO_PublicFpType : Type<And<[PPHLO_PublicType.predicate, 
+def PPHLO_PublicFpType : Type<And<[PPHLO_PublicType.predicate,
     CPred<"$_self.cast<::mlir::pphlo::PublicType>().getBase().isa<::mlir::FloatType>()">]>, "public fixed-point type", "::pphlo::PubFpType">;
 
-def PPHLO_SecretFpType : Type<And<[PPHLO_SecretType.predicate, 
+def PPHLO_SecretFpType : Type<And<[PPHLO_SecretType.predicate,
     CPred<"$_self.cast<::mlir::pphlo::SecretType>().getBase().isa<::mlir::FloatType>()">]>, "secret fixed-point type", "::pphlo::SecFpType">;
 
 def PPHLO_Tensor : StaticShapeTensorOf<[PPHLO_PublicType, PPHLO_SecretType]>;
diff --git a/libspu/kernel/hal/BUILD.bazel b/libspu/kernel/hal/BUILD.bazel
index 630b2d28..eead12ee 100644
--- a/libspu/kernel/hal/BUILD.bazel
+++ b/libspu/kernel/hal/BUILD.bazel
@@ -110,6 +110,7 @@ spu_cc_test(
     srcs = ["fxp_base_test.cc"],
     deps = [
         ":fxp_base",
+        ":type_cast",
         "//libspu/kernel:test_util",
     ],
 )
diff --git a/libspu/kernel/hal/README.md b/libspu/kernel/hal/README.md
index e37467b7..ab585b37 100644
--- a/libspu/kernel/hal/README.md
+++ b/libspu/kernel/hal/README.md
@@ -1,29 +1,39 @@
-# Object model
+# HAL
+
+## Object model
+
 Generally, we adopt concept-based polymorphism, it's widely used in MLIR, and it provides:
+
 1. runtime polymorphism, just like other object oriented system.
 2. value semantic, aka, RAII semantic.
 
 see [here](https://sean-parent.stlab.cc/papers-and-presentations/#value-semantics-and-concept-based-polymorphism) for more details.
 
-# Type system
+## Type system
+
 Type could be modeled as a tuple `<vistype, datatype, shape>`
+
 * **vistype**, the visibility type, could be one of {Public, Secret}.
 * **datatype**, the data encoding type, could be on of FXP, INT.
 * **shape**, the shape.
 
 Note
+
 * Share type (AShare/BShare/HShare) is not exposed to vm IR, so it's not exposed to IR type.
 * plaintext type (u8/i8/u32/i32/f32/f64/...) is 'readonly', we can only import/export them via `buffer view` objects.
 
-# Dispatching
+## Dispatching
+
 Due to the complexity of MPC value types, the dispatch chain is complicated.
 
 For example,
+
 ```python
 multiply :: (value<SFXP>, value<PINT>)         # multiply secret fxp to public integer.
 ```
 
 should be firstly dispatched via encoding type, result in:
+
 ```python
 multiply :: (SFXP, PINT)         # multiply secret fxp to public integer.
   multiply :: (SINT, PINT) -> SINT
@@ -31,6 +41,7 @@ multiply :: (SFXP, PINT)         # multiply secret fxp to public integer.
 ```
 
 then dispatch via visibility type, result in:
+
 ```python
 multiply :: (SFXP, PINT)         # multiply secret fxp to public integer.
   multiply :: (SINT, PINT) -> SINT
@@ -38,38 +49,44 @@ multiply :: (SFXP, PINT)         # multiply secret fxp to public integer.
   truncate :: (SINT)
 ```
 
-## Dispatching schema
+### Dispatching schema
+
 To model the dispatching rule, we use the following schema,
+
 1. make one dispatch do one thing, i.e. either lowering datatype or visibility type.
 2. make the dispatch layered, so they could be moved to compile-time progressively.
 
 Schema: `[<dtype>_]op_name[_<Visibility::type>]`
 
-- if `<dtype>` is missing, (op start with `_`), then dtype will be unchecked.
-- if `[<dtype>_]` not provided, then both fxp & int will be well handled.
-- if `[_<Visibility::type>]` not provided, then both secret/public will be well handled.
+* if `<dtype>` is missing, (op start with `_`), then dtype will be unchecked.
+* if `[<dtype>_]` not provided, then both fxp & int will be well handled.
+* if `[_<Visibility::type>]` not provided, then both secret/public will be well handled.
 
 For example:
-- `i_add_sp`: int add between secret and public.
--  `_add_sp`: unchecked add between secret & public, works on ring (integer)
--   `add_sp`: add between secret&public, both int, fxp will be handled.
-- `f_mul_pp`: fxp mul between public and public.
-- `f_mul   `: fxp mul, both secret&public will be handled.
+
+* `i_add_sp`: int add between secret and public.
+* `_add_sp`: unchecked add between secret & public, works on ring (integer)
+* `add_sp`: add between secret&public, both int, fxp will be handled.
+* `f_mul_pp`: fxp mul between public and public.
+* `f_mul`: fxp mul, both secret&public will be handled.
 
 If <dtype> is well specified, then output dtype will be certain, otherwise
 output dtype will set to CtType::UNKNOWN.
-- `i_add_sp`: out dtype is int.
--   `add_sp`: out dtype is specified, according to inputs type.
--  `_add_sp`: out dtype is UNKNOWN.
 
-## A complicated example
+* `i_add_sp`: out dtype is int.
+* `add_sp`: out dtype is specified, according to inputs type.
+* `_add_sp`: out dtype is UNKNOWN.
+
+### A complicated example
+
 `secret reciprocal` is a typical non-polynomial op, in this example, we use Newton-Raphson to find the approximated result.
 
 The approximation progress depends on other exp/square/add/mul/div, the tracing result is listed below.
 
-Note: `secret reciprocal` may dependends on `divide to public`, while general divide depends on reciprocal, to break the dependenty chain, the typed dispatch schema is import.
+Note: `secret reciprocal` may dependends on `divide to public`, while general divide depends on reciprocal, to break the dependenty chain,
+the typed dispatch schema is import.
 
-Please see [spu/modules/fxp.cc][] for details.
+Please see [spu/modules/fxp.cc] for details.
 
 ```python
 make_secret(BufferView)                    # Value a = make_secret(&ctx, 3.14f);
@@ -123,5 +140,8 @@ _s2p(value<SFXP>)
 dump_public(value<PFXP>)
 ```
 
-## Progressive lowering
-The lowering could be done in runtime (dynamic) or compile-time (static), the final goal is to make a thin-runtime (secret protocol only) and fat-compile-time (all data computation, numeric approximations..), but currently it's hard to move all into compile-time, so we layering the runtime dispatch rule and try to move them one by one to compile-time.
+### Progressive lowering
+
+The lowering could be done in runtime (dynamic) or compile-time (static), the final goal is to make a thin-runtime (secret protocol only) and
+fat-compile-time (all data computation, numeric approximations..), but currently it's hard to move all into compile-time, so we layering the
+runtime dispatch rule and try to move them one by one to compile-time.
diff --git a/libspu/kernel/hal/constants.cc b/libspu/kernel/hal/constants.cc
index f62d46d1..46353f22 100644
--- a/libspu/kernel/hal/constants.cc
+++ b/libspu/kernel/hal/constants.cc
@@ -46,19 +46,20 @@ Value make_pub2k(SPUContext* ctx, const PtBufferView& bv) {
 // TODO: formalize and test it.
 // clang-format off
 // NOLINTBEGIN, readability-implicit-bool-conversion, modernize-use-bool-literals
-bool kCastFlags[DT_FXP+1][DT_FXP+1] = { 
-//{_,  I1, I8, U8, I16,U16,I32,U32,I64,U64,FXP}
-  {0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0  },  // _
-  {0,  1,  1,  1,  1,  1,  1,  1,  1,  1,  0  },  // I1
-  {0,  0,  1,  0,  1,  0,  1,  0,  1,  0,  0  },  // I8
-  {0,  0,  1,  1,  1,  1,  1,  1,  1,  1,  0  },  // U8
-  {0,  0,  0,  0,  1,  0,  1,  0,  1,  0,  0  },  // I16
-  {0,  0,  0,  0,  1,  1,  1,  1,  1,  1,  0  },  // U16
-  {0,  0,  0,  0,  0,  0,  1,  0,  1,  0,  0  },  // I32
-  {0,  0,  0,  0,  0,  0,  1,  1,  1,  1,  0  },  // U32
-  {0,  0,  0,  0,  0,  0,  0,  0,  1,  0,  0  },  // I64
-  {0,  0,  0,  0,  0,  0,  0,  0,  1,  1,  0  },  // U64
-  {0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  1  },  // FXP
+bool kCastFlags[DT_F64+1][DT_F64+1] = {
+//{_,  I1, I8, U8, I16,U16,I32,U32,I64,U64,F32,F64}
+  {0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0  , 0},  // _
+  {0,  1,  1,  1,  1,  1,  1,  1,  1,  1,  0  , 0},  // I1
+  {0,  0,  1,  0,  1,  0,  1,  0,  1,  0,  0  , 0},  // I8
+  {0,  0,  1,  1,  1,  1,  1,  1,  1,  1,  0  , 0},  // U8
+  {0,  0,  0,  0,  1,  0,  1,  0,  1,  0,  0  , 0},  // I16
+  {0,  0,  0,  0,  1,  1,  1,  1,  1,  1,  0  , 0},  // U16
+  {0,  0,  0,  0,  0,  0,  1,  0,  1,  0,  0  , 0},  // I32
+  {0,  0,  0,  0,  0,  0,  1,  1,  1,  1,  0  , 0},  // U32
+  {0,  0,  0,  0,  0,  0,  0,  0,  1,  0,  0  , 0},  // I64
+  {0,  0,  0,  0,  0,  0,  0,  0,  1,  1,  0  , 0},  // U64
+  {0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  1  , 1},  // F32
+  {0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0  , 1},  // F64
 };
 // NOLINTEND
 // clang-format on
@@ -104,8 +105,8 @@ Value constant(SPUContext* ctx, PtBufferView init, 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);
+  if (dtype == DT_F32 || dtype == DT_F64) {
+    return constant(ctx, 0.0F, dtype, shape);
   } else {
     return constant(ctx, static_cast<uint8_t>(0), dtype, shape);
   }
@@ -119,8 +120,9 @@ Value iota(SPUContext* ctx, DataType dtype, int64_t numel) {
   });
 }
 
-Value epsilon(SPUContext* ctx, absl::Span<const int64_t> shape) {
-  return _constant(ctx, static_cast<int128_t>(1), shape).asFxp();
+Value epsilon(SPUContext* ctx, DataType dtype,
+              absl::Span<const int64_t> shape) {
+  return _constant(ctx, static_cast<int128_t>(1), shape).setDtype(dtype);
 }
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/constants.h b/libspu/kernel/hal/constants.h
index 9e202210..c456382c 100644
--- a/libspu/kernel/hal/constants.h
+++ b/libspu/kernel/hal/constants.h
@@ -41,6 +41,7 @@ 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(SPUContext* ctx, absl::Span<const int64_t> shape = {});
+Value epsilon(SPUContext* ctx, DataType dtype,
+              absl::Span<const int64_t> shape = {});
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/fxp_approx.cc b/libspu/kernel/hal/fxp_approx.cc
index 05269744..70e6d0bc 100644
--- a/libspu/kernel/hal/fxp_approx.cc
+++ b/libspu/kernel/hal/fxp_approx.cc
@@ -28,9 +28,9 @@ namespace spu::kernel::hal {
 namespace {
 
 // simple convenient function.
-Value f_constant(SPUContext* ctx, const PtBufferView& init,
+Value f_constant(SPUContext* ctx, const PtBufferView& init, DataType dtype,
                  absl::Span<int64_t const> shape) {
-  return constant(ctx, init, DT_FXP, shape);
+  return constant(ctx, init, dtype, shape);
 }
 
 }  // namespace
@@ -53,25 +53,25 @@ 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);
 
-  const auto p0 = f_constant(ctx, -0.205466671951 * 10, x.shape());
-  const auto p1 = f_constant(ctx, -0.88626599391 * 10, x.shape());
-  const auto p2 = f_constant(ctx, 0.610585199015 * 10, x.shape());
-  const auto p3 = f_constant(ctx, 0.481147460989 * 10, x.shape());
+  const auto p0 = f_constant(ctx, -0.205466671951F * 10, x.dtype(), x.shape());
+  const auto p1 = f_constant(ctx, -0.88626599391F * 10, x.dtype(), x.shape());
+  const auto p2 = f_constant(ctx, 0.610585199015F * 10, x.dtype(), x.shape());
+  const auto p3 = f_constant(ctx, 0.481147460989F * 10, x.dtype(), x.shape());
 
-  const auto q0 = f_constant(ctx, 0.353553425277, x.shape());
-  const auto q1 = f_constant(ctx, 0.454517087629 * 10, x.shape());
-  const auto q2 = f_constant(ctx, 0.642784209029 * 10, x.shape());
-  const auto q3 = f_constant(ctx, 0.1 * 10, x.shape());
+  const auto q0 = f_constant(ctx, 0.353553425277F, x.dtype(), x.shape());
+  const auto q1 = f_constant(ctx, 0.454517087629F * 10, x.dtype(), x.shape());
+  const auto q2 = f_constant(ctx, 0.642784209029F * 10, x.dtype(), x.shape());
+  const auto q3 = f_constant(ctx, 0.1F * 10, x.dtype(), x.shape());
 
   auto p2524 = _mul(ctx, x, p1);
   p2524 = _add(ctx, p2524, _mul(ctx, x2, p2));
   p2524 = _add(ctx, p2524, _mul(ctx, x3, p3));
-  p2524 = _add(ctx, _trunc(ctx, p2524), p0).asFxp();
+  p2524 = _add(ctx, _trunc(ctx, p2524), p0).setDtype(x.dtype());
 
   auto q2524 = _mul(ctx, x, q1);
   q2524 = _add(ctx, q2524, _mul(ctx, x2, q2));
   q2524 = _add(ctx, q2524, _mul(ctx, x3, q3));
-  q2524 = _add(ctx, _trunc(ctx, q2524), q0).asFxp();
+  q2524 = _add(ctx, _trunc(ctx, q2524), q0).setDtype(x.dtype());
 
   return detail::div_goldschmidt(ctx, p2524, q2524);
 }
@@ -90,7 +90,7 @@ Value log2_pade_approx(SPUContext* ctx, const Value& x) {
 
   // let x = x_norm * factor, where x in [0.5, 1.0)
   auto msb = detail::highestOneBit(ctx, x);
-  auto factor = _bitrev(ctx, msb, 0, 2 * num_fxp_bits).asFxp();
+  auto factor = _bitrev(ctx, msb, 0, 2 * num_fxp_bits).setDtype(x.dtype());
   detail::hintNumberOfBits(factor, 2 * num_fxp_bits);
   auto norm = f_mul(ctx, x, factor);
 
@@ -101,7 +101,7 @@ Value log2_pade_approx(SPUContext* ctx, const Value& x) {
              ctx, log2_pade_approx_for_normalized(ctx, norm),
              _lshift(ctx, _sub(ctx, k, _constant(ctx, num_fxp_bits, x.shape())),
                      num_fxp_bits))
-      .asFxp();
+      .setDtype(x.dtype());
 }
 
 // See P11, A.2.4 Logarithm and Exponent,
@@ -111,29 +111,31 @@ Value log2_pade_approx(SPUContext* ctx, const Value& x) {
 // Householder iterations. This approximation is accurate within 2% relative
 // error on [0.0001, 250].
 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_1 = f_div(ctx, x, f_constant(ctx, 120.0, x.dtype(), x.shape()));
   Value term_2 = f_mul(
       ctx,
       f_exp(ctx,
-            f_negate(ctx,
-                     f_add(ctx, f_mul(ctx, x, f_constant(ctx, 2.0, x.shape())),
-                           f_constant(ctx, 1.0, x.shape())))),
-      f_constant(ctx, 20.0, x.shape()));
-  Value y =
-      f_add(ctx, f_sub(ctx, term_1, term_2), f_constant(ctx, 3.0, x.shape()));
+            f_negate(
+                ctx,
+                f_add(ctx,
+                      f_mul(ctx, x, f_constant(ctx, 2.0, x.dtype(), x.shape())),
+                      f_constant(ctx, 1.0, x.dtype(), x.shape())))),
+      f_constant(ctx, 20.0, x.dtype(), x.shape()));
+  Value y = f_add(ctx, f_sub(ctx, term_1, term_2),
+                  f_constant(ctx, 3.0, x.dtype(), x.shape()));
 
   std::vector<Value> coeffs;
   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()));
+    coeffs.emplace_back(f_constant(ctx, 1.0 / (1.0 + i), x.dtype(), x.shape()));
   }
 
   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()),
+    Value h = f_sub(ctx, f_constant(ctx, 1.0, x.dtype(), x.shape()),
                     f_mul(ctx, x, f_exp(ctx, f_negate(ctx, y))));
     y = f_sub(ctx, y, detail::f_polynomial(ctx, h, coeffs));
   }
@@ -148,8 +150,8 @@ Value exp_taylor_series(SPUContext* ctx, const Value& x) {
   SPU_ENFORCE(fxp_exp_iters != 0, "fxp_exp_iters should not be {}",
               fxp_exp_iters);
 
-  Value res = f_add(ctx, _trunc(ctx, x, fxp_exp_iters).asFxp(),
-                    f_constant(ctx, 1.0, x.shape()));
+  Value res = f_add(ctx, _trunc(ctx, x, fxp_exp_iters).setDtype(x.dtype()),
+                    f_constant(ctx, 1.0F, x.dtype(), x.shape()));
 
   for (size_t i = 0; i < fxp_exp_iters; i++) {
     res = f_square(ctx, res);
@@ -172,12 +174,16 @@ Value exp2_pade_approx_for_positive_pure_decimal(SPUContext* ctx,
   auto x4 = f_mul(ctx, x, x3);
   auto x5 = f_mul(ctx, x, x4);
 
-  const auto p0 = f_constant(ctx, 0.100000007744302 * 10, x.shape());
-  const auto p1 = f_constant(ctx, 0.693147180426163, x.shape());
-  const auto p2 = f_constant(ctx, 0.240226510710170, x.shape());
-  const auto p3 = f_constant(ctx, 0.555040686204663 / 10, x.shape());
-  const auto p4 = f_constant(ctx, 0.961834122588046 / 100, x.shape());
-  const auto p5 = f_constant(ctx, 0.133273035928143 / 100, x.shape());
+  const auto p0 =
+      f_constant(ctx, 0.100000007744302F * 10, x.dtype(), x.shape());
+  const auto p1 = f_constant(ctx, 0.693147180426163F, x.dtype(), x.shape());
+  const auto p2 = f_constant(ctx, 0.240226510710170F, x.dtype(), x.shape());
+  const auto p3 =
+      f_constant(ctx, 0.555040686204663F / 10, x.dtype(), x.shape());
+  const auto p4 =
+      f_constant(ctx, 0.961834122588046F / 100, x.dtype(), x.shape());
+  const auto p5 =
+      f_constant(ctx, 0.133273035928143F / 100, x.dtype(), x.shape());
 
   auto res = _mul(ctx, x, p1);
   res = _add(ctx, res, _mul(ctx, x2, p2));
@@ -185,7 +191,7 @@ Value exp2_pade_approx_for_positive_pure_decimal(SPUContext* ctx,
   res = _add(ctx, res, _mul(ctx, x4, p4));
   res = _add(ctx, res, _mul(ctx, x5, p5));
 
-  return _add(ctx, _trunc(ctx, res), p0).asFxp();
+  return _add(ctx, _trunc(ctx, res), p0).setDtype(x.dtype());
 }
 
 // Refer to
@@ -204,7 +210,8 @@ Value exp2_pade_approx(SPUContext* ctx, const Value& x) {
   const auto x_bshare = _prefer_b(ctx, x);
   const auto x_msb = _rshift(ctx, x_bshare, bit_width - 1);
   auto x_integer = _rshift(ctx, x_bshare, fbits);
-  auto x_fraction = _sub(ctx, x, _lshift(ctx, x_integer, fbits)).asFxp();
+  auto x_fraction =
+      _sub(ctx, x, _lshift(ctx, x_integer, fbits)).setDtype(x.dtype());
   auto ret = exp2_pade_approx_for_positive_pure_decimal(ctx, x_fraction);
 
   for (size_t idx = 0; idx < int_bits; idx++) {
@@ -215,7 +222,7 @@ Value exp2_pade_approx(SPUContext* ctx, const Value& x) {
     ret = _mul(ctx, ret,
                _add(ctx, _mul(ctx, a, _constant(ctx, K, x.shape())),
                     _sub(ctx, k1, a)))
-              .asFxp();
+              .setDtype(ret.dtype());
   }
 
   // If we could ensure the integer bits of x is 5.
@@ -225,16 +232,19 @@ Value exp2_pade_approx(SPUContext* ctx, const Value& x) {
   //            (x_hat) - 32 = -x
   //  exp2(x_hat) / exp2(32) = exp(-x)
   //  so exp(-x) = exp2(x_hat) / exp2(32)
-  auto ret_reciprocal = _trunc(ctx, ret, std::pow(2, int_bits)).asFxp();
+  auto ret_reciprocal =
+      _trunc(ctx, ret, std::pow(2, int_bits)).setDtype(ret.dtype());
 
   // ret + msb * (reciprocal - ret)
-  return f_add(ctx, ret,
-               _mul(ctx, x_msb, f_sub(ctx, ret_reciprocal, ret)).asFxp());
+  return f_add(
+      ctx, ret,
+      _mul(ctx, x_msb, f_sub(ctx, ret_reciprocal, ret)).setDtype(ret.dtype()));
 }
 
 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())));
+  return f_exp2(ctx, f_mul(ctx, x,
+                           f_constant(ctx, std::log2(std::exp(1.0F)), x.dtype(),
+                                      x.shape())));
 }
 
 // Refer to
@@ -247,15 +257,16 @@ Value tanh_pade_approx(SPUContext* ctx, const Value& x) {
   const auto x_4 = f_square(ctx, x_2);
   const auto x_5 = f_mul(ctx, x_2, x_3);
 
-  const auto dividend =
-      f_add(ctx, x,
-            f_add(ctx, f_div(ctx, x_3, f_constant(ctx, 9.0, x.shape())),
-                  f_div(ctx, x_5, f_constant(ctx, 945.0, x.shape()))));
+  const auto dividend = f_add(
+      ctx, x,
+      f_add(ctx, f_div(ctx, x_3, f_constant(ctx, 9.0F, x.dtype(), x.shape())),
+            f_div(ctx, x_5, f_constant(ctx, 945.0F, x.dtype(), x.shape()))));
 
-  const auto divisor =
-      f_add(ctx, f_constant(ctx, 1.0, x.shape()),
-            f_add(ctx, f_div(ctx, x_2, f_constant(ctx, 9.0 / 4.0, x.shape())),
-                  f_div(ctx, x_4, f_constant(ctx, 63.0, x.shape()))));
+  const auto divisor = f_add(
+      ctx, f_constant(ctx, 1.0F, x.dtype(), x.shape()),
+      f_add(ctx,
+            f_div(ctx, x_2, f_constant(ctx, 9.0F / 4.0F, x.dtype(), x.shape())),
+            f_div(ctx, x_4, f_constant(ctx, 63.0F, x.dtype(), x.shape()))));
 
   return f_div(ctx, dividend, divisor);
 }
@@ -279,14 +290,14 @@ Value f_exp(SPUContext* ctx, const Value& x) {
       // The valid input for exp_pade_approx is [-kInputLimit, kInputLimit].
       // TODO(junfeng): should merge clamp into exp_pade_approx to save msb ops.
       const float kInputLimit = 32 / std::log2(std::exp(1));
-      const auto clampped_x =
-          _clamp(ctx, x, f_constant(ctx, -kInputLimit, x.shape()),
-                 f_constant(ctx, kInputLimit, x.shape()))
+      const auto clamped_x =
+          _clamp(ctx, x, f_constant(ctx, -kInputLimit, x.dtype(), x.shape()),
+                 f_constant(ctx, kInputLimit, x.dtype(), x.shape()))
               .setDtype(x.dtype());
-      return detail::exp_pade_approx(ctx, clampped_x);
+      return detail::exp_pade_approx(ctx, clamped_x);
     }
     default:
-      SPU_THROW("unexpected exp approxmation method {}",
+      SPU_THROW("unexpected exp approximation method {}",
                 ctx->config().fxp_exp_mode());
   }
 }
@@ -303,12 +314,12 @@ Value f_log(SPUContext* ctx, const Value& x) {
   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()),
+      return f_mul(ctx, f_constant(ctx, std::log(2.0F), x.dtype(), x.shape()),
                    f_log2(ctx, x));
     case RuntimeConfig::LOG_NEWTON:
       return detail::log_householder_approx(ctx, x);
     default:
-      SPU_THROW("unlogected log approxmation method {}",
+      SPU_THROW("unexpected log approximation method {}",
                 ctx->config().fxp_log_mode());
   }
 }
@@ -318,7 +329,7 @@ Value f_log1p(SPUContext* ctx, const Value& x) {
 
   SPU_ENFORCE(x.isFxp());
 
-  return f_log(ctx, f_add(ctx, f_constant(ctx, 1.0, x.shape()), x));
+  return f_log(ctx, f_add(ctx, f_constant(ctx, 1.0F, x.dtype(), x.shape()), x));
 }
 
 Value f_log2(SPUContext* ctx, const Value& x) {
@@ -326,7 +337,7 @@ Value f_log2(SPUContext* ctx, const Value& x) {
 
   SPU_ENFORCE(x.isFxp());
 
-  return detail::log2_pade_approx(ctx, x).asFxp();
+  return detail::log2_pade_approx(ctx, x).setDtype(x.dtype());
 }
 
 Value f_exp2(SPUContext* ctx, const Value& x) {
@@ -349,23 +360,25 @@ static Value rsqrt_init_guess(SPUContext* ctx, const Value& x, const Value& z) {
   auto z_rev = _bitrev(ctx, z, 0, 2 * f);
   detail::hintNumberOfBits(z_rev, 2 * f);
 
-  auto u = _trunc(ctx, _mul(ctx, x, z_rev)).asFxp();
+  auto u = _trunc(ctx, _mul(ctx, x, z_rev)).setDtype(x.dtype());
 
   // 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->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()),
-                                 f_constant(ctx, 26.02942339, x.shape())};
+    std::vector<Value> coeffs = {
+        f_constant(ctx, -15.47994394F, x.dtype(), x.shape()),
+        f_constant(ctx, 38.4714796F, x.dtype(), x.shape()),
+        f_constant(ctx, -49.86605845F, x.dtype(), x.shape()),
+        f_constant(ctx, 26.02942339F, x.dtype(), x.shape())};
     r = f_add(ctx, detail::f_polynomial(ctx, u, coeffs),
-              f_constant(ctx, 4.14285016, x.shape()));
+              f_constant(ctx, 4.14285016F, x.dtype(), x.shape()));
   } else {
-    std::vector<Value> coeffs = {f_constant(ctx, -5.9417, x.shape()),
-                                 f_constant(ctx, 4.7979, x.shape())};
+    std::vector<Value> coeffs = {
+        f_constant(ctx, -5.9417F, x.dtype(), x.shape()),
+        f_constant(ctx, 4.7979F, x.dtype(), x.shape())};
     r = f_add(ctx, detail::f_polynomial(ctx, u, coeffs),
-              f_constant(ctx, 3.1855, x.shape()));
+              f_constant(ctx, 3.1855F, x.dtype(), x.shape()));
   }
 
   return r;
@@ -449,11 +462,11 @@ Value f_rsqrt(SPUContext* ctx, const Value& x) {
     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();
+    return _trunc(ctx, _mul(ctx, r.get(), comp)).setDtype(x.dtype());
   } else {
     auto r = rsqrt_init_guess(ctx, x, z);
     auto comp = rsqrt_comp(ctx, x, z);
-    return _trunc(ctx, _mul(ctx, r, comp)).asFxp();
+    return _trunc(ctx, _mul(ctx, r, comp)).setDtype(x.dtype());
   }
 }
 
@@ -467,8 +480,8 @@ Value f_rsqrt(SPUContext* 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());
-  const auto c1 = f_constant(ctx, 1.5, x.shape());
+  const auto c0 = f_constant(ctx, 0.5F, x.dtype(), x.shape());
+  const auto c1 = f_constant(ctx, 1.5F, x.dtype(), x.shape());
 
   Value y0 = f_rsqrt(ctx, x);
   Value g = f_mul(ctx, x, y0);
@@ -490,14 +503,14 @@ namespace {
 
 Value sigmiod_real(SPUContext* ctx, const Value& x) {
   // f(x) = 1/(1+exp(-x))
-  const auto c1 = f_constant(ctx, 1.0F, x.shape());
+  const auto c1 = f_constant(ctx, 1.0F, x.dtype(), x.shape());
   return f_reciprocal(ctx, f_add(ctx, c1, f_exp(ctx, f_negate(ctx, 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());
+  const auto c1 = f_constant(ctx, 0.5F, x.dtype(), x.shape());
+  const auto c2 = f_constant(ctx, 0.125F, x.dtype(), x.shape());
   return f_add(ctx, c1, f_mul(ctx, c2, x));
 }
 
@@ -506,15 +519,15 @@ Value sigmiod_seg3(SPUContext* ctx, const Value& x) {
   //        1            if       x > 4
   //        0            if  -4 > x
   // Rounds = Gt + Mux*2 = 4 + Log(K)
-  auto upper = f_constant(ctx, 1.0F, x.shape());
-  auto lower = f_constant(ctx, 0.0F, x.shape());
+  auto upper = f_constant(ctx, 1.0F, x.dtype(), x.shape());
+  auto lower = f_constant(ctx, 0.0F, x.dtype(), x.shape());
   auto middle = sigmiod_mm1(ctx, x);
 
-  auto upper_bound = f_constant(ctx, 4.0F, x.shape());
-  auto lower_bound = f_constant(ctx, -4.0F, x.shape());
+  auto upper_bound = f_constant(ctx, 4.0F, x.dtype(), x.shape());
+  auto lower_bound = f_constant(ctx, -4.0F, x.dtype(), x.shape());
 
   auto ret = _mux(ctx, f_less(ctx, upper_bound, x), upper, middle);
-  return _mux(ctx, f_less(ctx, x, lower_bound), lower, ret).asFxp();
+  return _mux(ctx, f_less(ctx, x, lower_bound), lower, ret).setDtype(x.dtype());
 }
 
 }  // namespace
diff --git a/libspu/kernel/hal/fxp_base.cc b/libspu/kernel/hal/fxp_base.cc
index 1708992d..8b80e15a 100644
--- a/libspu/kernel/hal/fxp_base.cc
+++ b/libspu/kernel/hal/fxp_base.cc
@@ -49,7 +49,7 @@ Value f_polynomial(SPUContext* ctx, const Value& x,
     res = _add(ctx, res, _mul(ctx, x_pow, coeffs[i]));
   }
 
-  return _trunc(ctx, res).asFxp();
+  return _trunc(ctx, res).setDtype(x.dtype());
 }
 
 Value highestOneBit(SPUContext* ctx, const Value& x) {
@@ -102,13 +102,13 @@ Value div_goldschmidt(SPUContext* ctx, const Value& a, const Value& b) {
   auto is_negative = _msb(ctx, b);
   // insert ``prefer_a'' because the msb bit are used twice.
   is_negative = _prefer_a(ctx, is_negative);
-  auto b_abs = _mux(ctx, is_negative, _negate(ctx, b), b).asFxp();
+  auto b_abs = _mux(ctx, is_negative, _negate(ctx, b), b).setDtype(b.dtype());
 
   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}
   const size_t num_fxp_bits = ctx->getFxpBits();
-  auto factor = _bitrev(ctx, b_msb, 0, 2 * num_fxp_bits).asFxp();
+  auto factor = _bitrev(ctx, b_msb, 0, 2 * num_fxp_bits).setDtype(b.dtype());
   detail::hintNumberOfBits(factor, 2 * num_fxp_bits);
 
   // compute normalize x_abs, [0.5, 1)
@@ -117,11 +117,11 @@ Value div_goldschmidt(SPUContext* ctx, const Value& a, const Value& b) {
   // initial guess:
   //   w = 1/c ≈ 2.9142 - 2c when c >= 0.5 and c < 1
   const auto k2 = _constant(ctx, 2, c.shape());
-  const auto k2_9142 = constant(ctx, 2.9142F, DT_FXP, c.shape());
-  auto w = f_sub(ctx, k2_9142, _mul(ctx, k2, c).asFxp());
+  const auto k2_9142 = constant(ctx, 2.9142F, b.dtype(), c.shape());
+  auto w = f_sub(ctx, k2_9142, _mul(ctx, k2, c).setDtype(b.dtype()));
 
   // init r=w, e=1-c*w
-  const auto& k1_ = constant(ctx, 1.0F, DT_FXP, c.shape());
+  const auto& k1_ = constant(ctx, 1.0F, b.dtype(), c.shape());
   auto r = w;
   auto e = f_sub(ctx, k1_, f_mul_with_sign(ctx, c, w, SignType::POSITIVE));
 
@@ -141,7 +141,7 @@ Value div_goldschmidt(SPUContext* ctx, const Value& a, const Value& b) {
   // However, it might overflow when the input x is too small.
   r = f_mul(ctx, r, a);
   r = f_mul(ctx, r, factor);
-  return _mux(ctx, is_negative, _negate(ctx, r), r).asFxp();
+  return _mux(ctx, is_negative, _negate(ctx, r), r).setDtype(a.dtype());
 }
 
 Value reciprocal_goldschmidt_positive(SPUContext* ctx, const Value& b_abs) {
@@ -149,7 +149,8 @@ Value reciprocal_goldschmidt_positive(SPUContext* ctx, const Value& b_abs) {
 
   // factor = 2^{2f-m} = 2^{f-m} * 2^f, the fixed point repr of 2^{f-m}
   const size_t num_fxp_bits = ctx->getFxpBits();
-  auto factor = _bitrev(ctx, b_msb, 0, 2 * num_fxp_bits).asFxp();
+  auto factor =
+      _bitrev(ctx, b_msb, 0, 2 * num_fxp_bits).setDtype(b_abs.dtype());
   detail::hintNumberOfBits(factor, 2 * num_fxp_bits);
 
   // compute normalize x_abs, [0.5, 1)
@@ -158,11 +159,13 @@ Value reciprocal_goldschmidt_positive(SPUContext* ctx, const Value& b_abs) {
   // initial guess:
   //   w = 1/b = 2.9142 - 2c when c >= 0.5 and c < 1
   const auto k2 = _constant(ctx, 2, c.shape());
-  const auto k2_9142 = constant(ctx, 2.9142F, DT_FXP, c.shape());
-  auto w = f_mul(ctx, f_sub(ctx, k2_9142, _mul(ctx, k2, c).asFxp()), factor);
+  const auto k2_9142 = constant(ctx, 2.9142F, b_abs.dtype(), c.shape());
+  auto w =
+      f_mul(ctx, f_sub(ctx, k2_9142, _mul(ctx, k2, c).setDtype(b_abs.dtype())),
+            factor);
 
   // init r=a*w, e=1-b*w
-  const auto& k1_ = constant(ctx, 1.0F, DT_FXP, c.shape());
+  const auto& k1_ = constant(ctx, 1.0F, b_abs.dtype(), c.shape());
   auto r = w;
   auto e = f_sub(ctx, k1_, f_mul_with_sign(ctx, b_abs, w, SignType::POSITIVE));
 
@@ -189,9 +192,9 @@ Value reciprocal_goldschmidt(SPUContext* ctx, const Value& b) {
   auto is_negative = _msb(ctx, b);
   is_negative = _prefer_a(ctx, is_negative);
 
-  auto b_abs = _mux(ctx, is_negative, _negate(ctx, b), b).asFxp();
+  auto b_abs = _mux(ctx, is_negative, _negate(ctx, b), b).setDtype(b.dtype());
   auto r = reciprocal_goldschmidt_positive(ctx, b_abs);
-  return _mux(ctx, is_negative, _negate(ctx, r), r).asFxp();
+  return _mux(ctx, is_negative, _negate(ctx, r), r).setDtype(b.dtype());
 }
 
 }  // namespace detail
@@ -200,7 +203,7 @@ Value f_negate(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_LEAF(ctx, x);
 
   SPU_ENFORCE(x.isFxp());
-  return _negate(ctx, x).asFxp();
+  return _negate(ctx, x).setDtype(x.dtype());
 }
 
 Value f_abs(SPUContext* ctx, const Value& x) {
@@ -209,7 +212,7 @@ Value f_abs(SPUContext* ctx, const Value& x) {
   SPU_ENFORCE(x.isFxp());
   const Value sign = _sign(ctx, x);
 
-  return _mul(ctx, sign, x).asFxp();
+  return _mul(ctx, sign, x).setDtype(x.dtype());
 }
 
 Value f_reciprocal(SPUContext* ctx, const Value& x) {
@@ -226,45 +229,43 @@ Value f_reciprocal(SPUContext* ctx, const Value& x) {
 Value f_add(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
-  SPU_ENFORCE(x.isFxp());
-  SPU_ENFORCE(y.isFxp());
+  SPU_ENFORCE(x.isFxp() && y.isFxp() && x.dtype() == y.dtype());
 
-  return _add(ctx, x, y).asFxp();
+  return _add(ctx, x, y).setDtype(x.dtype());
 }
 
 Value f_sub(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
-  SPU_ENFORCE(x.isFxp());
-  SPU_ENFORCE(y.isFxp());
+  SPU_ENFORCE(x.isFxp() && y.isFxp() && x.dtype() == y.dtype());
+
   return f_add(ctx, x, f_negate(ctx, y));
 }
 
 Value f_mul(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
-  SPU_ENFORCE(x.isFxp());
-  SPU_ENFORCE(y.isFxp());
-  return _trunc(ctx, _mul(ctx, x, y)).asFxp();
+  SPU_ENFORCE(x.isFxp() && y.isFxp() && x.dtype() == y.dtype());
+
+  return _trunc(ctx, _mul(ctx, x, y)).setDtype(x.dtype());
 }
 
 Value f_mul_with_sign(SPUContext* ctx, const Value& x, const Value& y,
                       SignType sign) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
-  SPU_ENFORCE(x.isFxp());
-  SPU_ENFORCE(y.isFxp());
+  SPU_ENFORCE(x.isFxp() && y.isFxp() && x.dtype() == y.dtype());
 
   switch (sign) {
     case SignType::POSITIVE:
       return _trunc_with_sign(ctx, _mul(ctx, x, y), ctx->getFxpBits(),
                               /*positive*/ true)
-          .asFxp();
+          .setDtype(x.dtype());
       break;
     case SignType::NEGATIVE:
       return _trunc_with_sign(ctx, _mul(ctx, x, y), ctx->getFxpBits(),
                               /*positive*/ false)
-          .asFxp();
+          .setDtype(x.dtype());
       break;
     case SignType::UNKNOWN:
     default:
@@ -276,10 +277,9 @@ Value f_mul_with_sign(SPUContext* 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());
-  SPU_ENFORCE(y.isFxp());
+  SPU_ENFORCE(x.isFxp() && y.isFxp() && x.dtype() == y.dtype());
 
-  return _trunc(ctx, _mmul(ctx, x, y)).asFxp();
+  return _trunc(ctx, _mmul(ctx, x, y)).setDtype(x.dtype());
 }
 
 Value f_conv2d(SPUContext* ctx, const Value& x, const Value& y,
@@ -287,17 +287,16 @@ Value f_conv2d(SPUContext* ctx, const Value& x, const Value& y,
                absl::Span<const int64_t> result_shape) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
-  SPU_ENFORCE(x.isFxp());
-  SPU_ENFORCE(y.isFxp());
+  SPU_ENFORCE(x.isFxp() && y.isFxp() && x.dtype() == y.dtype());
 
-  return _trunc(ctx, _conv2d(ctx, x, y, window_strides, result_shape)).asFxp();
+  return _trunc(ctx, _conv2d(ctx, x, y, window_strides, result_shape))
+      .setDtype(x.dtype());
 }
 
 Value f_div(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
 
-  SPU_ENFORCE(x.isFxp());
-  SPU_ENFORCE(y.isFxp());
+  SPU_ENFORCE(x.isFxp() && y.isFxp() && x.dtype() == y.dtype());
 
   if (x.isPublic() && y.isPublic()) {
     return f_div_p(ctx, x, y);
@@ -309,8 +308,7 @@ Value f_div(SPUContext* 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());
-  SPU_ENFORCE(y.isFxp());
+  SPU_ENFORCE(x.isFxp() && y.isFxp() && x.dtype() == y.dtype());
 
   return _equal(ctx, x, y).setDtype(DT_I1);
 }
@@ -318,8 +316,7 @@ Value f_equal(SPUContext* 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());
-  SPU_ENFORCE(y.isFxp());
+  SPU_ENFORCE(x.isFxp() && y.isFxp() && x.dtype() == y.dtype());
 
   return _less(ctx, x, y).setDtype(DT_I1);
 }
@@ -332,7 +329,7 @@ Value f_square(SPUContext* ctx, const Value& x) {
   // TODO(juhou) can use truncate with msb=0
 
   return _trunc_with_sign(ctx, _mul(ctx, x, x), ctx->getFxpBits(), true)
-      .asFxp();
+      .setDtype(x.dtype());
 }
 
 Value f_floor(SPUContext* ctx, const Value& x) {
@@ -341,7 +338,7 @@ Value f_floor(SPUContext* ctx, const Value& x) {
   SPU_ENFORCE(x.isFxp());
 
   const size_t fbits = ctx->getFxpBits();
-  return _lshift(ctx, _arshift(ctx, x, fbits), fbits).asFxp();
+  return _lshift(ctx, _arshift(ctx, x, fbits), fbits).setDtype(x.dtype());
 }
 
 Value f_ceil(SPUContext* ctx, const Value& x) {
@@ -350,8 +347,9 @@ Value f_ceil(SPUContext* ctx, const Value& x) {
   SPU_ENFORCE(x.isFxp());
 
   // ceil(x) = floor(x + 1.0 - epsilon)
-  const auto k1 = constant(ctx, 1.0F, DT_FXP, x.shape());
-  return f_floor(ctx, f_add(ctx, x, f_sub(ctx, k1, epsilon(ctx, x.shape()))));
+  const auto k1 = constant(ctx, 1.0F, x.dtype(), x.shape());
+  return f_floor(
+      ctx, f_add(ctx, x, f_sub(ctx, k1, epsilon(ctx, x.dtype(), x.shape()))));
 }
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/fxp_cleartext.cc b/libspu/kernel/hal/fxp_cleartext.cc
index e3c2fa0e..bf8e9cbd 100644
--- a/libspu/kernel/hal/fxp_cleartext.cc
+++ b/libspu/kernel/hal/fxp_cleartext.cc
@@ -25,7 +25,7 @@ template <typename 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());
+  SPU_ENFORCE(in.isFxp(), "expected fxp, got={}", in.dtype());
 
   const size_t fxp_bits = ctx->getFxpBits();
   const auto field = in.storage_type().as<Ring2k>()->field();
@@ -42,7 +42,7 @@ Value applyFloatingPointFn(SPUContext* ctx, const Value& in, FN&& fn) {
   DataType dtype;
   const auto out = encodeToRing(f32_arr, field, fxp_bits, &dtype);
 
-  SPU_ENFORCE(dtype == DT_FXP, "sanity failed");
+  SPU_ENFORCE(dtype == DT_F32 || dtype == DT_F64, "sanity failed");
   return Value(out.as(in.storage_type()), dtype);
 }
 
@@ -52,8 +52,8 @@ Value applyFloatingPointFn(SPUContext* ctx, const Value& x, const Value& y,
   SPU_TRACE_HAL_DISP(ctx, x, y);
   SPU_ENFORCE(x.isPublic() && y.isPublic(), "expect public, got {}, {}",
               x.vtype(), y.vtype());
-  SPU_ENFORCE((x.dtype() == DT_FXP) && (y.dtype() == DT_FXP),
-              "expected fxp, got={} {}", x.dtype(), y.dtype());
+  SPU_ENFORCE((x.isFxp()) && (y.isFxp()), "expected fxp, got={} {}", x.dtype(),
+              y.dtype());
   SPU_ENFORCE(x.shape() == y.shape());
 
   const auto field = x.storage_type().as<Ring2k>()->field();
@@ -73,7 +73,7 @@ Value applyFloatingPointFn(SPUContext* ctx, const Value& x, const Value& y,
 
   DataType dtype;
   const auto out = encodeToRing(flp_x, field, fxp_bits, &dtype);
-  SPU_ENFORCE(dtype == DT_FXP, "sanity failed");
+  SPU_ENFORCE(dtype == DT_F32 || dtype == DT_F64, "sanity failed");
   return Value(out.as(x.storage_type()), dtype);
 }
 
diff --git a/libspu/kernel/hal/polymorphic.cc b/libspu/kernel/hal/polymorphic.cc
index a2990d9f..761a2bf8 100644
--- a/libspu/kernel/hal/polymorphic.cc
+++ b/libspu/kernel/hal/polymorphic.cc
@@ -51,9 +51,9 @@ Value dtypeBinaryDispatch(std::string_view op_name, SPUContext* ctx,
     return FnInt(ctx, dtype_cast(ctx, x, common_type),
                  dtype_cast(ctx, y, common_type));
   } else if (x.isInt() && y.isFxp()) {
-    return FnFxp(ctx, dtype_cast(ctx, x, DT_FXP), y);
+    return FnFxp(ctx, dtype_cast(ctx, x, y.dtype()), y);
   } else if (x.isFxp() && y.isInt()) {
-    return FnFxp(ctx, x, dtype_cast(ctx, y, DT_FXP));
+    return FnFxp(ctx, x, dtype_cast(ctx, y, x.dtype()));
   } else if (x.isFxp() && y.isFxp()) {
     return FnFxp(ctx, x, y);
   } else {
@@ -92,12 +92,14 @@ Value sub(SPUContext* 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();
+  auto new_dtype = x.isFxp() ? x.dtype() : y.dtype();
+  return _mul(ctx, x, y).setDtype(new_dtype);
 }
 
 Value mixed_mmul(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_LEAF(ctx, x, y);
-  return _mmul(ctx, x, y).asFxp();
+  auto new_dtype = x.isFxp() ? x.dtype() : y.dtype();
+  return _mmul(ctx, x, y).setDtype(new_dtype);
 }
 
 Value mul(SPUContext* ctx, const Value& x, const Value& y) {
@@ -192,7 +194,9 @@ Value abs(SPUContext* ctx, const Value& x) {
 Value exp(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
-  return f_exp(ctx, dtype_cast(ctx, in, DT_FXP));
+  SPU_ENFORCE(in.isFxp());
+
+  return f_exp(ctx, in);
 }
 
 Value select(SPUContext* ctx, const Value& pred, const Value& a,
@@ -244,19 +248,24 @@ Value logistic(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
   SPU_ENFORCE(in.isFxp());
+
   return f_sigmoid(ctx, in);
 }
 
 Value log(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
-  return f_log(ctx, dtype_cast(ctx, in, DT_FXP));
+  SPU_ENFORCE(in.isFxp());
+
+  return f_log(ctx, in);
 }
 
 Value log1p(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
-  return f_log1p(ctx, dtype_cast(ctx, in, DT_FXP));
+  SPU_ENFORCE(in.isFxp());
+
+  return f_log1p(ctx, in);
 }
 
 Value reciprocal(SPUContext* ctx, const Value& in) {
@@ -300,6 +309,14 @@ Value min(SPUContext* ctx, const Value& x, const Value& y) {
 
 Value power(SPUContext* ctx, const Value& x, const Value& y) {
   SPU_TRACE_HAL_DISP(ctx, x, y);
+
+  if (x.isInt() && y.isInt()) {
+    auto x_f = dtype_cast(ctx, x, DT_F32);
+    auto y_f = dtype_cast(ctx, y, DT_F32);
+    auto ret = power(ctx, x_f, y_f);
+    return dtype_cast(ctx, ret, x.dtype());
+  }
+
   // x^y = e^(y*ln(x))
   return exp(ctx, mul(ctx, y, log(ctx, x)));
 }
@@ -313,8 +330,8 @@ Value idiv(SPUContext* ctx, const Value& x, const Value& y) {
 
   Value q;
   {
-    const auto x_f = dtype_cast(ctx, abs_x, DT_FXP);
-    const auto y_f = dtype_cast(ctx, abs_y, DT_FXP);
+    const auto x_f = dtype_cast(ctx, abs_x, DT_F32);
+    const auto y_f = dtype_cast(ctx, abs_y, DT_F32);
 
     auto approx_q = div(ctx, x_f, y_f);
 
@@ -342,8 +359,11 @@ Value div(SPUContext* ctx, const Value& x, const Value& y) {
     return idiv(ctx, x, y);
   }
 
-  const auto x_f = dtype_cast(ctx, x, DT_FXP);
-  const auto y_f = dtype_cast(ctx, y, DT_FXP);
+  // Kind of a hack to compute max dtype
+  auto dtype = std::max(x.dtype(), y.dtype());
+
+  const auto x_f = dtype_cast(ctx, x, dtype);
+  const auto y_f = dtype_cast(ctx, y, dtype);
 
 #define F_DIV_WITH_DIRECT_GOLDSCHMIDT_METHOD
 #ifdef F_DIV_WITH_DIRECT_GOLDSCHMIDT_METHOD
@@ -394,24 +414,29 @@ Value right_shift_arithmetic(SPUContext* ctx, const Value& x, size_t bits) {
 Value log2(SPUContext* ctx, const Value& in) {
   SPU_TRACE_HAL_DISP(ctx, in);
 
-  return f_log2(ctx, dtype_cast(ctx, in, DT_FXP));
+  SPU_ENFORCE(in.isFxp());
+
+  return f_log2(ctx, in);
 }
 
 Value exp2(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
-  return f_exp2(ctx, dtype_cast(ctx, x, DT_FXP));
+  SPU_ENFORCE(x.isFxp());
+
+  return f_exp2(ctx, x);
 }
 
 Value tanh(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
+  SPU_ENFORCE(x.isFxp());
+
   // For tanh inputs beyond [-3, 3], result is infinitely close to -1, 1
   // pade approximation has a relative ok result between [-3, 3], so clamp
   // inputs to this range.
-  auto normalized_x = clamp(ctx, dtype_cast(ctx, x, DT_FXP),
-                            constant(ctx, -3.F, DT_FXP, x.shape()),
-                            constant(ctx, 3.F, DT_FXP, x.shape()));
+  auto normalized_x = clamp(ctx, x, constant(ctx, -3.F, x.dtype(), x.shape()),
+                            constant(ctx, 3.F, x.dtype(), x.shape()));
 
   return f_tanh(ctx, normalized_x);
 }
@@ -419,13 +444,17 @@ Value tanh(SPUContext* 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));
+  SPU_ENFORCE(x.isFxp());
+
+  return f_rsqrt(ctx, x);
 }
 
 Value sqrt(SPUContext* ctx, const Value& x) {
   SPU_TRACE_HAL_DISP(ctx, x);
 
-  return f_sqrt(ctx, dtype_cast(ctx, x, DT_FXP));
+  SPU_ENFORCE(x.isFxp());
+
+  return f_sqrt(ctx, x);
 }
 
 Value sign(SPUContext* ctx, const Value& x) {
diff --git a/libspu/kernel/hal/prot_wrapper.h b/libspu/kernel/hal/prot_wrapper.h
index be890cc3..3e858add 100644
--- a/libspu/kernel/hal/prot_wrapper.h
+++ b/libspu/kernel/hal/prot_wrapper.h
@@ -14,6 +14,8 @@
 
 #pragma once
 
+#include <optional>
+
 #include "libspu/core/context.h"
 #include "libspu/core/value.h"
 
diff --git a/libspu/kernel/hal/random.cc b/libspu/kernel/hal/random.cc
index 56ad1f48..d16537b8 100644
--- a/libspu/kernel/hal/random.cc
+++ b/libspu/kernel/hal/random.cc
@@ -53,7 +53,7 @@ Value rng_uniform(SPUContext* ctx, const Value& lo, const Value& hi,
     ele = dist(gen);
   }
 
-  return constant(ctx, buffer, DT_FXP);
+  return constant(ctx, buffer, lo.dtype());
 }
 
 Value random(SPUContext* ctx, Visibility vis, DataType dtype,
diff --git a/libspu/kernel/hal/type_cast.cc b/libspu/kernel/hal/type_cast.cc
index 172463d2..b48667d7 100644
--- a/libspu/kernel/hal/type_cast.cc
+++ b/libspu/kernel/hal/type_cast.cc
@@ -22,11 +22,11 @@
 namespace spu::kernel::hal {
 namespace {
 
-Value int2fxp(SPUContext* ctx, const Value& x) {
+Value int2fxp(SPUContext* ctx, const Value& x, DataType to_type) {
   SPU_TRACE_HAL_LEAF(ctx, x);
   SPU_ENFORCE(x.isInt(), "expect integer, got {}", x.dtype());
 
-  return _lshift(ctx, x, ctx->getFxpBits()).asFxp();
+  return _lshift(ctx, x, ctx->getFxpBits()).setDtype(to_type);
 }
 
 // Casting fxp to integer.
@@ -46,7 +46,7 @@ Value int2fxp(SPUContext* ctx, const Value& x) {
 //
 Value fxp2int(SPUContext* ctx, const Value& x, DataType to_type) {
   SPU_TRACE_HAL_LEAF(ctx, x);
-  SPU_ENFORCE(x.dtype() == DataType::DT_FXP);
+  SPU_ENFORCE(x.isFxp());
 
   const size_t fxp_bits = ctx->getFxpBits();
   const Value kOneMinusEps = _constant(ctx, (1 << fxp_bits) - 1, x.shape());
@@ -86,16 +86,16 @@ Value dtype_cast(SPUContext* ctx, const Value& in, DataType to_type) {
       return Value(in.data(), to_type);
     } else {
       SPU_ENFORCE(isFixedPoint(to_type));
-      return int2fxp(ctx, in);
+      return int2fxp(ctx, in, to_type);
     }
   } else {
     if (isInteger(to_type)) {
       return fxp2int(ctx, in, to_type);
     } else {
-      SPU_ENFORCE(to_type == DT_FXP, "expect to_type FXP, got {}", to_type);
-      SPU_ENFORCE(in.dtype() == DT_FXP, "expect in type FXP, got {}", to_type);
-      // we only support one FXP type, do nothing.
-      return in;
+      SPU_ENFORCE(to_type == DT_F32 || to_type == DT_F64,
+                  "expect to_type FXP, got {}", to_type);
+      SPU_ENFORCE(in.isFxp(), "expect in type FXP, got {}", in.dtype());
+      return Value(in.data(), to_type);
     }
   }
 
diff --git a/libspu/kernel/hlo/basic_binary.cc b/libspu/kernel/hlo/basic_binary.cc
index edb5ac9c..5c5f791e 100644
--- a/libspu/kernel/hlo/basic_binary.cc
+++ b/libspu/kernel/hlo/basic_binary.cc
@@ -15,6 +15,7 @@
 #include "libspu/kernel/hlo/basic_binary.h"
 
 #include "libspu/kernel/hal/constants.h"
+#include "libspu/kernel/hal/debug.h"
 #include "libspu/kernel/hal/polymorphic.h"
 #include "libspu/kernel/hal/type_cast.h"
 
diff --git a/libspu/kernel/hlo/basic_unary.cc b/libspu/kernel/hlo/basic_unary.cc
index c74efe95..c973883a 100644
--- a/libspu/kernel/hlo/basic_unary.cc
+++ b/libspu/kernel/hlo/basic_unary.cc
@@ -46,8 +46,8 @@ 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.
-  return hal::sub(ctx, hal::exp(ctx, in),
-                  hal::constant(ctx, 1.0, DT_FXP, in.shape()));
+  auto e = hal::exp(ctx, in);
+  return hal::sub(ctx, e, hal::constant(ctx, 1.0F, e.dtype(), in.shape()));
 }
 
 spu::Value Not(SPUContext *ctx, const spu::Value &in) {
@@ -77,7 +77,7 @@ spu::Value Round_AFZ(SPUContext *ctx, const spu::Value &in) {
   SPU_ENFORCE(in.isFxp(), "Round only supports fxp");
 
   auto sign_in = hal::sign(ctx, in);
-  auto p_half = hal::constant(ctx, 0.5, DT_FXP, in.shape());
+  auto p_half = hal::constant(ctx, 0.5F, in.dtype(), in.shape());
   p_half = hal::mul(ctx, sign_in, p_half);
 
   auto round = hal::add(ctx, in, p_half);
diff --git a/libspu/kernel/hlo/const.cc b/libspu/kernel/hlo/const.cc
index 6713539f..4a6dadeb 100644
--- a/libspu/kernel/hlo/const.cc
+++ b/libspu/kernel/hlo/const.cc
@@ -36,6 +36,8 @@ spu::Value Iota(SPUContext *ctx, DataType dtype, int64_t numel) {
   return hal::iota(ctx, dtype, numel);
 }
 
-spu::Value Epsilon(SPUContext *ctx) { return hal::epsilon(ctx); }
+spu::Value Epsilon(SPUContext *ctx, DataType dtype) {
+  return hal::epsilon(ctx, dtype);
+}
 
 }  // namespace spu::kernel::hlo
diff --git a/libspu/kernel/hlo/const.h b/libspu/kernel/hlo/const.h
index 951b2ca5..7f2098db 100644
--- a/libspu/kernel/hlo/const.h
+++ b/libspu/kernel/hlo/const.h
@@ -30,6 +30,6 @@ 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(SPUContext *ctx);
+spu::Value Epsilon(SPUContext *ctx, DataType dtype);
 
 }  // namespace spu::kernel::hlo
diff --git a/libspu/mpc/ab_api.cc b/libspu/mpc/ab_api.cc
index 9a400e19..b5d75966 100644
--- a/libspu/mpc/ab_api.cc
+++ b/libspu/mpc/ab_api.cc
@@ -397,7 +397,7 @@ Value add_bb(SPUContext* ctx, const Value& x, const Value& y) {
   }
 }
 
-Value carry_out(SPUContext* ctx, const Value& x, const Value& y, size_t k) {
+Value carry_a2b(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);
diff --git a/libspu/mpc/ab_api.h b/libspu/mpc/ab_api.h
index 78b6cadc..45c3fb90 100644
--- a/libspu/mpc/ab_api.h
+++ b/libspu/mpc/ab_api.h
@@ -103,6 +103,6 @@ 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);
+Value carry_a2b(SPUContext* ctx, const Value& x, const Value& y, size_t k);
 
 }  // namespace spu::mpc
diff --git a/libspu/mpc/aby3/conversion.cc b/libspu/mpc/aby3/conversion.cc
index dc42a833..e4a60ba1 100644
--- a/libspu/mpc/aby3/conversion.cc
+++ b/libspu/mpc/aby3/conversion.cc
@@ -568,7 +568,7 @@ ArrayRef MsbA2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
   auto wrap_m = WrapValue(m, shape);
   auto wrap_n = WrapValue(n, shape);
   {
-    auto carry = carry_out(sctx, wrap_m, wrap_n, nbits);
+    auto carry = carry_a2b(sctx, wrap_m, wrap_n, nbits);
 
     // Compute the k'th bit.
     //   (m^n)[k] ^ carry
diff --git a/libspu/mpc/aby3/conversion.h b/libspu/mpc/aby3/conversion.h
index 4f146097..c3a4a27c 100644
--- a/libspu/mpc/aby3/conversion.h
+++ b/libspu/mpc/aby3/conversion.h
@@ -105,13 +105,13 @@ class MsbA2B : public UnaryKernel {
   static constexpr char kBindName[] = "msb_a2b";
 
   ce::CExpr latency() const override {
-    // 1 * carry_out : log(k) + 1
+    // 1 * carry : log(k) + 1
     // 1 * rotate: 1
     return Log(ce::K()) + 1 + 1;
   }
 
   ce::CExpr comm() const override {
-    // 1 * carry_out: k + 2 * k
+    // 1 * carry : k + 2 * k
     // 1 * rotate: k
     return ce::K() * 4;
   }
diff --git a/libspu/mpc/cheetah/arith/cheetah_dot.cc b/libspu/mpc/cheetah/arith/cheetah_dot.cc
index 4af0fde2..b1422f96 100644
--- a/libspu/mpc/cheetah/arith/cheetah_dot.cc
+++ b/libspu/mpc/cheetah/arith/cheetah_dot.cc
@@ -738,13 +738,6 @@ CheetahDot::CheetahDot(std::shared_ptr<yacl::link::Context> lctx) {
 
 CheetahDot::~CheetahDot() = default;
 
-CheetahDot::CheetahDot(std::unique_ptr<Impl> impl) : impl_(std::move(impl)) {}
-
-std::unique_ptr<CheetahDot> CheetahDot::Fork() {
-  auto ptr = new CheetahDot(impl_->Fork());
-  return std::unique_ptr<CheetahDot>(ptr);
-}
-
 ArrayRef CheetahDot::DotOLE(const ArrayRef &inp, yacl::link::Context *conn,
                             const Shape3D &dim3, bool is_lhs) {
   SPU_ENFORCE(impl_ != nullptr);
diff --git a/libspu/mpc/cheetah/arith/cheetah_dot.h b/libspu/mpc/cheetah/arith/cheetah_dot.h
index a064c205..cefac197 100644
--- a/libspu/mpc/cheetah/arith/cheetah_dot.h
+++ b/libspu/mpc/cheetah/arith/cheetah_dot.h
@@ -33,8 +33,6 @@ class CheetahDot {
 
   ~CheetahDot();
 
-  std::unique_ptr<CheetahDot> Fork();
-
   CheetahDot& operator=(const CheetahDot&) = delete;
 
   CheetahDot(const CheetahDot&) = delete;
@@ -62,8 +60,6 @@ class CheetahDot {
  private:
   struct Impl;
 
-  explicit CheetahDot(std::unique_ptr<Impl> impl);
-
   std::unique_ptr<Impl> impl_{nullptr};
 
   std::shared_ptr<yacl::link::Context> lctx_{nullptr};
diff --git a/libspu/mpc/cheetah/arith/cheetah_mul.cc b/libspu/mpc/cheetah/arith/cheetah_mul.cc
index 6b306a59..826d2a29 100644
--- a/libspu/mpc/cheetah/arith/cheetah_mul.cc
+++ b/libspu/mpc/cheetah/arith/cheetah_mul.cc
@@ -55,6 +55,7 @@ struct CheetahMul::Impl : public EnableCPRNG {
   static constexpr int kNoiseFloodRandomBits = 50;
 
   static constexpr size_t kParallelGrain = 1;
+  static constexpr size_t kCtAsyncParallel = 8;
 
   explicit Impl(std::shared_ptr<yacl::link::Context> lctx)
       : lctx_(std::move(lctx)) {
@@ -81,28 +82,6 @@ struct CheetahMul::Impl : public EnableCPRNG {
     return parms;
   }
 
-  std::unique_ptr<Impl> Fork() {
-    using namespace seal;
-    auto f = std::make_unique<Impl>(lctx_->Spawn());
-    std::unique_lock guard(context_lock_);
-
-    f->current_crt_plain_bitlen_ = current_crt_plain_bitlen_;
-    if (seal_cntxts_.empty()) {
-      return f;
-    }
-
-    f->seal_cntxts_ = seal_cntxts_;
-    f->ms_helpers_ = ms_helpers_;
-    // NOTE(juhou): just take the pointer
-    f->secret_key_ = secret_key_;
-    f->pair_public_key_ = pair_public_key_;
-    f->sym_encryptors_ = sym_encryptors_;
-    f->decryptors_ = decryptors_;
-    f->pk_encryptors_ = pk_encryptors_;
-    f->bfv_encoders_ = bfv_encoders_;
-    return f;
-  }
-
   size_t num_slots() const { return parms_.poly_modulus_degree(); }
 
   void LazyExpandSEALContexts(uint32_t field_bitlen,
@@ -228,7 +207,7 @@ struct CheetahMul::Impl : public EnableCPRNG {
 };
 
 void CheetahMul::Impl::LazyInitModSwitchHelper(uint32_t field_bitlen) {
-  // TODO(juhou): multi-thread safe for ModulusSwitchHelper ?
+  std::unique_lock guard(context_lock_);
   if (ms_helpers_.count(field_bitlen) > 0) {
     return;
   }
@@ -251,6 +230,7 @@ void CheetahMul::Impl::LazyInitModSwitchHelper(uint32_t field_bitlen) {
   parms.set_coeff_modulus(crt_modulus);
 
   seal::SEALContext crt_context(parms, false, seal::sec_level_type::none);
+  SPU_ENFORCE(crt_context.parameters_set());
   ms_helpers_.emplace(field_bitlen,
                       ModulusSwitchHelper(crt_context, field_bitlen));
 }
@@ -300,6 +280,11 @@ void CheetahMul::Impl::LazyExpandSEALContexts(uint32_t field_bitlen,
 
   auto crt_modulus =
       seal::CoeffModulus::Create(parms_.poly_modulus_degree(), crt_moduli_bits);
+  // NOTE(juhou): sort the primes to make sure new primes are placed in the back
+  std::sort(crt_modulus.begin(), crt_modulus.end(),
+            [](const seal::Modulus &p, const seal::Modulus &q) {
+              return p.value() > q.value();
+            });
   uint32_t current_num_ctx = seal_cntxts_.size();
 
   for (uint32_t i = current_num_ctx; i < num_seal_ctx; ++i) {
@@ -446,9 +431,14 @@ size_t CheetahMul::Impl::EncryptArrayThenSend(const ArrayRef &array,
   if (conn == nullptr) {
     conn = lctx_.get();
   }
+
   int nxt_rank = conn->NextRank();
-  for (auto &ct : payload) {
-    conn->SendAsync(nxt_rank, ct, "");
+  for (size_t i = 0; i < payload.size(); i += kCtAsyncParallel) {
+    size_t this_batch = std::min(payload.size() - i, kCtAsyncParallel);
+    conn->Send(nxt_rank, payload[i], "");
+    for (size_t j = 1; j < this_batch; ++j) {
+      conn->SendAsync(nxt_rank, payload[i + j], "");
+    }
   }
   return payload.size();
 }
@@ -598,9 +588,14 @@ ArrayRef CheetahMul::Impl::MuThenResponse(
   if (conn == nullptr) {
     conn = lctx_.get();
   }
+
   int nxt_rank = conn->NextRank();
-  for (auto &ct : response) {
-    conn->SendAsync(nxt_rank, ct, fmt::format("Send to P{}", nxt_rank));
+  for (size_t i = 0; i < response.size(); i += kCtAsyncParallel) {
+    size_t this_batch = std::min(response.size() - i, kCtAsyncParallel);
+    conn->Send(nxt_rank, response[i], "");
+    for (size_t j = 1; j < this_batch; ++j) {
+      conn->SendAsync(nxt_rank, response[i + j], "");
+    }
   }
 
   for (auto &pt : ecd_random) {
@@ -724,17 +719,10 @@ CheetahMul::CheetahMul(std::shared_ptr<yacl::link::Context> lctx) {
   impl_ = std::make_unique<Impl>(lctx);
 }
 
-CheetahMul::CheetahMul(std::unique_ptr<Impl> impl) { impl_ = std::move(impl); }
-
 CheetahMul::~CheetahMul() = default;
 
 int CheetahMul::Rank() const { return impl_->Rank(); }
 
-std::unique_ptr<CheetahMul> CheetahMul::Fork() {
-  auto ptr = new CheetahMul(impl_->Fork());
-  return std::unique_ptr<CheetahMul>(ptr);
-}
-
 size_t CheetahMul::OLEBatchSize() const {
   SPU_ENFORCE(impl_ != nullptr);
   return impl_->OLEBatchSize();
diff --git a/libspu/mpc/cheetah/arith/cheetah_mul.h b/libspu/mpc/cheetah/arith/cheetah_mul.h
index 238a2431..bb24f815 100644
--- a/libspu/mpc/cheetah/arith/cheetah_mul.h
+++ b/libspu/mpc/cheetah/arith/cheetah_mul.h
@@ -32,8 +32,6 @@ class CheetahMul {
 
   ~CheetahMul();
 
-  std::unique_ptr<CheetahMul> Fork();
-
   CheetahMul& operator=(const CheetahMul&) = delete;
 
   CheetahMul(const CheetahMul&) = delete;
@@ -52,9 +50,6 @@ class CheetahMul {
  private:
   struct Impl;
 
-  // Used by Fork
-  explicit CheetahMul(std::unique_ptr<Impl> impl);
-
   std::unique_ptr<Impl> impl_{nullptr};
 };
 
diff --git a/libspu/mpc/cheetah/arith/vector_encoder.h b/libspu/mpc/cheetah/arith/vector_encoder.h
index a510836b..68b4120b 100644
--- a/libspu/mpc/cheetah/arith/vector_encoder.h
+++ b/libspu/mpc/cheetah/arith/vector_encoder.h
@@ -29,7 +29,7 @@ class VectorEncoder {
   // clang-format off
   // Math:
   //    Forward(x) * Backward(y) mod X^N + 1 gives the inner product <x, y>
-  // When doing in encryption, we need the encrypted part to be scaled up by some fixed `Delta` 
+  // When doing in encryption, we need the encrypted part to be scaled up by some fixed `Delta`
   // For example Enc(Forward(Delta*x)) * Backward(y) can give the <x,  y> without error.
   // Or we can encrypt the backward part, i.e., Enc(Backward(Delta*y)) * Forward(x) also gives the <x, y> without error.
   // clang-format on
@@ -39,6 +39,8 @@ class VectorEncoder {
 
   const ModulusSwitchHelper &ms_helper() const { return *ms_helper_; }
 
+  size_t poly_degree() const { return poly_deg_; }
+
  private:
   size_t poly_deg_{0};
   std::shared_ptr<ModulusSwitchHelper> ms_helper_;
diff --git a/libspu/mpc/cheetah/ot/basic_ot_prot.cc b/libspu/mpc/cheetah/ot/basic_ot_prot.cc
index e4c5041c..c3a72ea2 100644
--- a/libspu/mpc/cheetah/ot/basic_ot_prot.cc
+++ b/libspu/mpc/cheetah/ot/basic_ot_prot.cc
@@ -35,12 +35,6 @@ BasicOTProtocols::BasicOTProtocols(std::shared_ptr<Communicator> conn)
   }
 }
 
-std::unique_ptr<BasicOTProtocols> BasicOTProtocols::Fork() {
-  // TODO(juhou) we can take from cached ROTs from the caller
-  auto conn = std::make_shared<Communicator>(conn_->lctx()->Spawn());
-  return std::make_unique<BasicOTProtocols>(conn);
-}
-
 BasicOTProtocols::~BasicOTProtocols() { Flush(); }
 
 void BasicOTProtocols::Flush() {
diff --git a/libspu/mpc/cheetah/ot/basic_ot_prot.h b/libspu/mpc/cheetah/ot/basic_ot_prot.h
index 5175a03e..d4522511 100644
--- a/libspu/mpc/cheetah/ot/basic_ot_prot.h
+++ b/libspu/mpc/cheetah/ot/basic_ot_prot.h
@@ -25,8 +25,6 @@ class BasicOTProtocols {
 
   ~BasicOTProtocols();
 
-  std::unique_ptr<BasicOTProtocols> Fork();
-
   int Rank() const;
 
   ArrayRef B2A(const ArrayRef &inp);
diff --git a/libspu/mpc/cheetah/state.h b/libspu/mpc/cheetah/state.h
index ecfb3ed4..b5899f5d 100644
--- a/libspu/mpc/cheetah/state.h
+++ b/libspu/mpc/cheetah/state.h
@@ -51,13 +51,6 @@ class CheetahMulState : public State {
 
   CheetahMul* get() { return mul_prot_.get(); }
 
-  bool hasLowCostFork() const override { return true; }
-
-  std::unique_ptr<State> fork() override {
-    auto ptr = new CheetahMulState(mul_prot_->Fork());
-    return std::unique_ptr<State>(ptr);
-  }
-
   std::array<ArrayRef, 3> TakeCachedBeaver(FieldType field, int64_t num);
 };
 
@@ -78,13 +71,6 @@ class CheetahDotState : public State {
   ~CheetahDotState() override = default;
 
   CheetahDot* get() { return dot_prot_.get(); }
-
-  bool hasLowCostFork() const override { return true; }
-
-  std::unique_ptr<State> fork() override {
-    auto ptr = new CheetahDotState(dot_prot_->Fork());
-    return std::unique_ptr<State>(ptr);
-  }
 };
 
 class CheetahOTState : public State {
diff --git a/libspu/mpc/common/communicator.h b/libspu/mpc/common/communicator.h
index d92ea12d..781d3931 100644
--- a/libspu/mpc/common/communicator.h
+++ b/libspu/mpc/common/communicator.h
@@ -23,6 +23,7 @@
 #include "yacl/base/buffer.h"
 #include "yacl/link/link.h"
 
+#include "libspu/core/array_ref.h"
 #include "libspu/core/object.h"
 #include "libspu/core/parallel_utils.h"
 #include "libspu/core/prelude.h"
diff --git a/libspu/mpc/semi2k/conversion.cc b/libspu/mpc/semi2k/conversion.cc
index 956e2ba1..1d0cdcdc 100644
--- a/libspu/mpc/semi2k/conversion.cc
+++ b/libspu/mpc/semi2k/conversion.cc
@@ -171,7 +171,7 @@ ArrayRef MsbA2B::proc(KernelEvalContext* ctx, const ArrayRef& in) const {
   auto m = WrapValue(bshrs[0], shape);
   auto n = WrapValue(bshrs[1], shape);
   {
-    auto carry = carry_out(sctx, m, n, k);
+    auto carry = carry_a2b(sctx, m, n, k);
 
     // Compute the k'th bit.
     //   (m^n)[k] ^ carry
diff --git a/libspu/mpc/semi2k/conversion.h b/libspu/mpc/semi2k/conversion.h
index 01a12263..0234f141 100644
--- a/libspu/mpc/semi2k/conversion.h
+++ b/libspu/mpc/semi2k/conversion.h
@@ -79,12 +79,12 @@ class MsbA2B : public UnaryKernel {
   static constexpr char kBindName[] = "msb_a2b";
 
   ce::CExpr latency() const override {
-    // 1 * carry_out : log(k) + 1
+    // 1 * carry: log(k) + 1
     return Log(ce::K()) + 1;
   }
 
   ce::CExpr comm() const override {
-    // 1 * carry_out: k + 4 * k
+    // 1 * carry: k + 4 * k
     return ce::K() * 5;
   }
 
diff --git a/libspu/mpc/tools/README.md b/libspu/mpc/tools/README.md
index 704de2ab..8cd7de25 100644
--- a/libspu/mpc/tools/README.md
+++ b/libspu/mpc/tools/README.md
@@ -4,7 +4,7 @@
 
 You can get help info like this:
 
-```
+```sh
 $ bazel run -c opt libspu/mpc/tools/benchmark -- --help
 USAGE: benchmark [options]
 
@@ -36,14 +36,14 @@ Generic Options:
 Because many parameters have default parameters, we don't have to specify each parameter.
 If you want run it with **standalone mode**, you can do this:
 
-```bash
+```sh
 bazel run -c opt libspu/mpc/tools/benchmark -- --benchmark_counters_tabular=true
 ```
 
 If you want **mparty mode** on localhost, you need start multi processes to simulate different parties, and we only care output of rank 0.
 eg: run **aby3** **mparty** benchmark as follows, you can create a script.:
 
-```
+```sh
 $ cat run_mpart_bench.sh
 #!/bin/sh
 bazel run -c opt //libspu/mpc/tools/benchmark -- --mode=mparty --rank=1 2>&1 >/dev/null &
@@ -54,7 +54,7 @@ $ sh run_mpart_bench.sh
 
 or
 
-```
+```sh
 sh docs/reference/run_benchmark.sh --output LAN.json --mode mparty
 ```
 
@@ -62,14 +62,14 @@ sh docs/reference/run_benchmark.sh --output LAN.json --mode mparty
 
 You can use **--benchmark_out=*.json --benchmark_out_format=json** to specify the output json or **docs/reference/run_benchmark.sh**, eg:
 
-```
+```sh
 $ bazel run -c opt //libspu/mpc/tools:benchmark -- --mode=standalone \
   --benchmark_out=standalone.json --benchmark_out_format=json
 ```
 
 or
 
-```
+```sh
 sh docs/reference/run_benchmark.sh --output standalone.json
 ```
 
@@ -80,7 +80,7 @@ To set network limitation manually, execute following commands on each docker or
 then specify ip with `parties` parameter, eg:
 set 300Mbps bandwidth，add 20ms latency:
 
-```
+```sh
 $ tc qdisc del dev eth0 root
 $ tc qdisc add dev eth0 root handle 1: tbf rate 300mbit burst 256kb latency 800ms
 $ tc qdisc add dev eth0 parent 1:1 handle 10: netem delay 20msec limit 8000
@@ -96,7 +96,7 @@ You can always specify output json filename with **--output**.
 
 ### - standalone
 
-```
+```sh
 sh docs/reference/run_benchmark.sh
 ```
 
@@ -107,7 +107,7 @@ This will generate standalone.json.
 For example, you can run this on 172.17.0.8, and use 172.17.0.1 as rank 1&2.
 This outputs WAN.json. Before run this script, you need **ssh-copy-id** to remote machine.
 
-```
+```sh
 sh docs/reference/run_benchmark.sh --mode mparty --remote huocun@172.17.0.1 --remote_dir /home/ssd0/huocun/ppu --remote_python_env /home/ssd0/huocun/ppu/venv/bin/activate --parties 172.17.0.8:9444,172.17.0.1:9445,172.17.0.1:9446 --output LAN.json
 ```
 
@@ -118,7 +118,7 @@ New round-trip delay will be old round-trip delay add delay values you set.
 Tool tc needs to be installed. (You can run `yum install -y iproute-tc` to install tc.)
 This will generate WAN_300mbit_20msec.json.
 
-```
+```sh
 sh docs/reference/run_benchmark.sh --mode mparty --remote huocun@172.17.0.1 --remote_dir /home/ssd0/huocun/ppu --remote_python_env /home/ssd0/huocun/ppu/venv/bin/activate --parties 172.17.0.8:9444,172.17.0.1:9445,172.17.0.1:9446 --rate 300mbit --delay 20msec
 ```
 
@@ -137,7 +137,8 @@ python docs/reference/gen_benchmark_report.py --output=report.md
     --sheet="Benchmark Protocol"
 ```
 
-This command can generate a xlsx report(/csv/md/html is also supported), rows will be grouped by **op_name** and **field_type**, columns with **env**, **buf_len**.
+This command can generate a xlsx report(/csv/md/html is also supported), rows will be grouped by **op_name** and **field_type**,
+columns with **env**, **buf_len**.
 **env** is from **input** filenames, add different protocol will be in different sheets when we specify
 **--output=*.md**.
 **time** is formated by reasonable unit, raw data is "real_time".
diff --git a/libspu/psi/core/labeled_psi/README.md b/libspu/psi/core/labeled_psi/README.md
index 653776c6..a70a8362 100644
--- a/libspu/psi/core/labeled_psi/README.md
+++ b/libspu/psi/core/labeled_psi/README.md
@@ -1,6 +1,9 @@
-# Introduction
+# LabelPSI
 
-Fully homomorphic encryption (FHE) can be used to build efficient (labeled) Private Set Intersection protocols in the unbalanced setting, where one of the sets is much larger than the other.
+## Introduction
+
+Fully homomorphic encryption (FHE) can be used to build efficient (labeled) Private Set Intersection protocols in the unbalanced setting,
+where one of the sets is much larger than the other.
 
 Reference:
 
@@ -10,15 +13,17 @@ Reference:
 - [FHE and Private Set Intersection](https://simons.berkeley.edu/talks/fhe-and-private-set-intersection)
 - [Private set intersection via somewhat homomorphic encryption by Ilia Iliashenko](https://fhe.org/talks/private-set-intersection)
 
-Microsoft [APSI (Asymmetric PSI) library](https://github.com/microsoft/APSI) provides a PSI functionality for asymmetric set sizes based on the latest [Labeled PSI protocol](https://eprint.iacr.org/2021/1116).
+Microsoft [APSI (Asymmetric PSI) library](https://github.com/microsoft/APSI) provides a PSI functionality for asymmetric set sizes based
+on the latest [Labeled PSI protocol](https://eprint.iacr.org/2021/1116).
+
+This is a wrap of [APSI library](https://github.com/microsoft/APSI), can be used for
 
-This is a wrap of [APSI library](https://github.com/microsoft/APSI), can be used for 
 - Unbalanced PSI
 - Malicious PSI
 - Labeled PSI
 - Keyword PIR
 
-# LabelPSI Protocol dataflow
+## LabelPSI Protocol dataflow
 
 |        | Client(Receiver) |         | Server(Sender)  |
 | ------ | ---------------- | ------- | --------------- |
@@ -30,7 +35,7 @@ This is a wrap of [APSI library](https://github.com/microsoft/APSI), can be used
 | Step 4 | Request Query    | ------> |                 |
 |        |                  | <------ | Response Query  |
 
-# LabelPSI source code
+## LabelPSI source code
 
 |     | file            | class              | function            |
 | --- | --------------- | ------------------ | ------------------- |
@@ -55,7 +60,7 @@ This is a wrap of [APSI library](https://github.com/microsoft/APSI), can be used
 |     |                 |                    | GetBinBundleCount   |
 |     |                 |                    | GetPackingRate      |
 
-# LabelPSI Parameters
+## LabelPSI Parameters
 
 |     | file        | class                | function                                                            |
 | --- | ----------- | -------------------- | ------------------------------------------------------------------- |
@@ -83,8 +88,7 @@ This is a wrap of [APSI library](https://github.com/microsoft/APSI), can be used
 
 We select three SEALParams from [APSI parameters](https://github.com/microsoft/APSI/tree/main/parameters) for different receiver and sender items size.
 
-# Security Tips
+## Security Tips
 
 Warning:  Labeled PSI are malicious PSI protocols, but malicious sender may do attack with his CuckooHash Table.
 We suggest use Labeled PSI protocol as one-way PSI, i.e., just Client(Receiver) gets the final intersection result.
-
diff --git a/libspu/spu.proto b/libspu/spu.proto
index 3d584257..a9435e70 100644
--- a/libspu/spu.proto
+++ b/libspu/spu.proto
@@ -37,7 +37,8 @@ enum DataType {
   DT_U32 = 7;   // uint32
   DT_I64 = 8;   // int64
   DT_U64 = 9;   // uint64
-  DT_FXP = 10;  // Fixed point type.
+  DT_F32 = 10;  // float
+  DT_F64 = 11;  // double
 }
 
 // The visibility type.
@@ -319,28 +320,38 @@ message CompilationSource {
   repeated Visibility input_visibility = 3;
 }
 
+enum XLAPrettyPrintKind {
+  TEXT = 0;
+  DOT = 1;
+  HTML = 2;
+}
+
 message CompilerOptions {
   // Pretty print
   bool enable_pretty_print = 1;
   string pretty_print_dump_dir = 2;
+  XLAPrettyPrintKind xla_pp_kind = 3;
   
   // Disable sqrt(x) + eps to sqrt(x+eps) rewrite
-  bool disable_sqrt_plus_epsilon_rewrite = 3;
+  bool disable_sqrt_plus_epsilon_rewrite = 10;
 
   // Disable x/sqrt(y) to x*rsqrt(y) rewrite
-  bool disable_div_sqrt_rewrite = 4;
+  bool disable_div_sqrt_rewrite = 11;
 
   // Disable reduce truncation optimization
-  bool disable_reduce_truncation_optimization = 5;
+  bool disable_reduce_truncation_optimization = 12;
 
   // Disable maxpooling optimization
-  bool disable_maxpooling_optimization = 6;
+  bool disable_maxpooling_optimization = 13;
 
   // Disallow mix type operations
-  bool disallow_mix_types_opts = 7;
+  bool disallow_mix_types_opts = 14;
 
   // Disable SelectOp optimization
-  bool disable_select_optimization = 8;
+  bool disable_select_optimization = 15;
+
+  // Enable optimize x/bcast(y) -> x * bcast(1/y)
+  bool enable_optimize_denominator_with_broadcast = 16;
 }
 
 // The executable format accepted by SPU runtime.
diff --git a/pyproject.toml b/pyproject.toml
index 1711619a..d613f1f8 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -3,3 +3,18 @@ skip-string-normalization = true
 
 [tool.isort]
 profile = "black"
+
+[tool.rstcheck]
+report_level = "ERROR"
+ignore_directives = [
+    "include",
+    "mermaid",
+    "autoclass",
+    "autofunction",
+]
+ignore_roles = [
+    "spu_code_host",
+]
+ignore_languages = [
+    "cpp"
+]
diff --git a/spu/utils/distributed.py b/spu/utils/distributed.py
index 811d4c0e..34a2b13d 100644
--- a/spu/utils/distributed.py
+++ b/spu/utils/distributed.py
@@ -371,7 +371,8 @@ def shape_spu_to_np(spu_shape):
 
 def dtype_spu_to_np(spu_dtype):
     MAP = {
-        spu_pb2.DataType.DT_FXP: np.float32,
+        spu_pb2.DataType.DT_F32: np.float32,
+        spu_pb2.DataType.DT_F64: np.float64,
         spu_pb2.DataType.DT_I1: np.bool_,
         spu_pb2.DataType.DT_I8: np.int8,
         spu_pb2.DataType.DT_U8: np.uint8,
diff --git a/spu/utils/frontend.py b/spu/utils/frontend.py
index d810f292..067229de 100644
--- a/spu/utils/frontend.py
+++ b/spu/utils/frontend.py
@@ -96,8 +96,8 @@ class Kind(Enum):
 def compile(
     kind: Kind,
     fn: Callable,
-    args: List,
-    kwargs: Dict,
+    m_args: List,
+    m_kwargs: Dict,
     input_names: List[str],
     input_vis: List,
     outputNameGen: Callable,
@@ -109,7 +109,7 @@ def compile(
 
         patches = _patch_jax()
 
-        ir_text, output = _jax_compilation(fn, static_argnums, args, kwargs)
+        ir_text, output = _jax_compilation(fn, static_argnums, m_args, m_kwargs)
 
         _restore_jax_patch(patches)
 
@@ -120,11 +120,11 @@ def compile(
         import tensorflow as tf
 
         tf_fn = tf.function(fn, jit_compile=True, experimental_relax_shapes=True)
-        ir_text = tf_fn.experimental_get_compiler_ir(*args, **kwargs)(
+        ir_text = tf_fn.experimental_get_compiler_ir(*m_args, **m_kwargs)(
             stage="hlo_serialized",
         )
 
-        cf = tf_fn.get_concrete_function(*args, **kwargs)
+        cf = tf_fn.get_concrete_function(*m_args, **m_kwargs)
 
         # TODO(junfeng): support input captures.
         assert (
@@ -144,7 +144,7 @@ def compile(
         ), "currently only torch.nn.Module is supported"
 
         # convert numpy.ndarray to torch tensor as torch_mlir required
-        arg_tensors = [torch.Tensor(arg) for arg in args]
+        arg_tensors = [torch.Tensor(arg) for arg in m_args]
         # get mlir module
         module = torch_mlir.compile(
             fn, arg_tensors, output_type=torch_mlir.OutputType.MHLO
diff --git a/spu/version.py b/spu/version.py
index 49ab6fed..f9eb4db9 100644
--- a/spu/version.py
+++ b/spu/version.py
@@ -13,4 +13,4 @@
 # limitations under the License.
 
 
-__version__ = "0.4.0b0"
+__version__ = "0.4.0b1"


Test patch:
diff --git a/libspu/compiler/tests/decompose_comparison.mlir b/libspu/compiler/tests/decompose_comparison.mlir
index f8bfdb33..c88d0626 100644
--- a/libspu/compiler/tests/decompose_comparison.mlir
+++ b/libspu/compiler/tests/decompose_comparison.mlir
@@ -42,4 +42,3 @@ func.func @main(%arg0: tensor<2x2x!pphlo.pub<i32>>, %arg1: tensor<2x2x!pphlo.pub
     %5 = "pphlo.greater_equal"(%arg0, %arg1) : (tensor<2x2x!pphlo.pub<i32>>, tensor<2x2x!pphlo.pub<i32>>) -> tensor<2x2x!pphlo.pub<i1>>
     return %0, %1, %2, %3, %4, %5 : tensor<2x2x!pphlo.pub<i1>>, tensor<2x2x!pphlo.pub<i1>>, tensor<2x2x!pphlo.pub<i1>>, tensor<2x2x!pphlo.pub<i1>>, tensor<2x2x!pphlo.pub<i1>>, tensor<2x2x!pphlo.pub<i1>>
   }
-  
\ No newline at end of file
diff --git a/libspu/compiler/tests/enum_conversion_test.cc b/libspu/compiler/tests/enum_conversion_test.cc
index aa9018ee..3d98fde3 100644
--- a/libspu/compiler/tests/enum_conversion_test.cc
+++ b/libspu/compiler/tests/enum_conversion_test.cc
@@ -35,10 +35,10 @@ TEST(EnumConversion, Public) {
 }
 
 TEST(EnumConversion, ProtoKinds) {
-  Visibility v;
+  Visibility v = Visibility::VIS_PUBLIC;
   // This is basically a compile time check...that there is no new field
   switch (v) {
-  case Visibility::VIS_SECRET:
+  case Visibility::VIS_SECRET: // NOLINT
     break;
   case Visibility::VIS_PUBLIC:
     break;
@@ -52,10 +52,10 @@ TEST(EnumConversion, ProtoKinds) {
 }
 
 TEST(EnumConversion, TdKinds) {
-  mlir::pphlo::Visibility v;
+  mlir::pphlo::Visibility v = mlir::pphlo::Visibility::VIS_PUBLIC;
   // This is basically a compile time check...that there is no new field
   switch (v) {
-  case mlir::pphlo::Visibility::VIS_SECRET:
+  case mlir::pphlo::Visibility::VIS_SECRET: // NOLINT
     break;
   case mlir::pphlo::Visibility::VIS_PUBLIC:
     break;
diff --git a/libspu/compiler/tests/hlo_to_pphlo_ops_nullary.mlir b/libspu/compiler/tests/hlo_to_pphlo_ops_nullary.mlir
index 8d62e8c3..1e6792d9 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_ops_nullary.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_ops_nullary.mlir
@@ -3,25 +3,25 @@
 func.func @main() -> (tensor<2x2xi1>) {
     // CHECK: "pphlo.constant"() {value = dense<true> : tensor<2x2xi1>} : () -> tensor<2x2x!pphlo.pub<i1>>
     %0 = stablehlo.constant dense<true> : tensor<2x2xi1>
-    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xi8>} : () -> tensor<2x2x!pphlo.pub<i8>> 
+    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xi8>} : () -> tensor<2x2x!pphlo.pub<i8>>
     %1 = stablehlo.constant dense<1> : tensor<2x2xi8>
-    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xi16>} : () -> tensor<2x2x!pphlo.pub<i16>> 
+    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xi16>} : () -> tensor<2x2x!pphlo.pub<i16>>
     %2 = stablehlo.constant dense<1> : tensor<2x2xi16>
-    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xi32>} : () -> tensor<2x2x!pphlo.pub<i32>> 
+    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xi32>} : () -> tensor<2x2x!pphlo.pub<i32>>
     %3 = stablehlo.constant dense<1> : tensor<2x2xi32>
-    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xi64>} : () -> tensor<2x2x!pphlo.pub<i64>> 
+    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xi64>} : () -> tensor<2x2x!pphlo.pub<i64>>
     %4 = stablehlo.constant dense<1> : tensor<2x2xi64>
-    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui8>} : () -> tensor<2x2x!pphlo.pub<ui8>> 
+    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui8>} : () -> tensor<2x2x!pphlo.pub<ui8>>
     %5 = stablehlo.constant dense<1> : tensor<2x2xui8>
-    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui16>} : () -> tensor<2x2x!pphlo.pub<ui16>> 
+    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui16>} : () -> tensor<2x2x!pphlo.pub<ui16>>
     %6 = stablehlo.constant dense<1> : tensor<2x2xui16>
-    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui32>} : () -> tensor<2x2x!pphlo.pub<ui32>> 
+    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui32>} : () -> tensor<2x2x!pphlo.pub<ui32>>
     %7 = stablehlo.constant dense<1> : tensor<2x2xui32>
-    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui64>} : () -> tensor<2x2x!pphlo.pub<ui64>> 
+    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui64>} : () -> tensor<2x2x!pphlo.pub<ui64>>
     %8 = stablehlo.constant dense<1> : tensor<2x2xui64>
-    // CHECK: "pphlo.constant"() {value = dense<1.000000e+00> : tensor<2x2xf32>} : () -> tensor<2x2x!pphlo.pub<f32>> 
+    // CHECK: "pphlo.constant"() {value = dense<1.000000e+00> : tensor<2x2xf32>} : () -> tensor<2x2x!pphlo.pub<f32>>
     %9 = stablehlo.constant dense<1.0> : tensor<2x2xf32>
-    // CHECK: "pphlo.constant"() {value = dense<1.000000e+00> : tensor<2x2xf64>} : () -> tensor<2x2x!pphlo.pub<f64>> 
+    // CHECK: "pphlo.constant"() {value = dense<1.000000e+00> : tensor<2x2xf64>} : () -> tensor<2x2x!pphlo.pub<f64>>
     %10 = stablehlo.constant dense<1.0> : tensor<2x2xf64>
     return %0 : tensor<2x2xi1>
 }
diff --git a/libspu/compiler/tests/hlo_to_pphlo_ops_other.mlir b/libspu/compiler/tests/hlo_to_pphlo_ops_other.mlir
index 878e6448..218d39ee 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_ops_other.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_ops_other.mlir
@@ -2,13 +2,13 @@
 
 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>>
-    %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = dense<1> : tensor<1xi64>} : (tensor<16xf32>) -> tensor<1024x16xf32>    
+    %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = dense<1> : tensor<1xi64>} : (tensor<16xf32>) -> tensor<1024x16xf32>
     // CHECK: %1 = "pphlo.reshape"(%arg0) : (tensor<16x!pphlo.pub<f32>>) -> tensor<1x16x!pphlo.pub<f32>>
     %1 = "stablehlo.reshape"(%arg0) : (tensor<16xf32>) -> tensor<1x16xf32>
     // CHECK: %2 = "pphlo.transpose"(%1) {permutation = dense<[1, 0]> : tensor<2xi64>} : (tensor<1x16x!pphlo.pub<f32>>) -> tensor<16x1x!pphlo.pub<f32>>
     %2 = "stablehlo.transpose"(%1) {permutation = dense<[1, 0]> : tensor<2xi64>} : (tensor<1x16xf32>) -> tensor<16x1xf32>
     // CHECK: %3 = "pphlo.dot"(%0, %2) : (tensor<1024x16x!pphlo.pub<f32>>, tensor<16x1x!pphlo.pub<f32>>) -> tensor<1024x1x!pphlo.pub<f32>>
-    %3 = "stablehlo.dot"(%0, %2) {precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]} : (tensor<1024x16xf32>, tensor<16x1xf32>) -> tensor<1024x1xf32> 
+    %3 = "stablehlo.dot"(%0, %2) {precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]} : (tensor<1024x16xf32>, tensor<16x1xf32>) -> tensor<1024x1xf32>
     // CHECK: %6 = "pphlo.concatenate"(%4, %5) {dimension = 1 : i64} : (tensor<1024x16x!pphlo.pub<f32>>, tensor<1024x1x!pphlo.pub<f32>>) -> tensor<1024x17x!pphlo.pub<f32>>
     %4 = "stablehlo.concatenate"(%0, %3) {dimension = 1 : i64} : (tensor<1024x16xf32>, tensor<1024x1xf32>) -> tensor<1024x17xf32>
     // CHECK: %7 = "pphlo.select"(%arg1, %arg2, %arg3) : (tensor<1024x1x!pphlo.pub<i1>>, tensor<1024x1x!pphlo.pub<f32>>, tensor<1024x1x!pphlo.pub<f32>>) -> tensor<1024x1x!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 e60e4bde..2fbaea36 100644
--- a/libspu/compiler/tests/hlo_to_pphlo_select_and_scatter.mlir
+++ b/libspu/compiler/tests/hlo_to_pphlo_select_and_scatter.mlir
@@ -18,6 +18,6 @@ func.func @main(%arg0: tensor<128x5x5x32xf32>, %arg1: tensor<128x4x4x32xf32>, %a
     ^bb0(%arg3: tensor<f32>, %arg4: tensor<f32>):
         %1 = stablehlo.add %arg3, %arg4 : tensor<f32>
         "stablehlo.return"(%1) : (tensor<f32>) -> ()
-    }) {padding = dense<0> : tensor<4x2xi64>, window_dimensions = dense<[1, 2, 2, 1]> : tensor<4xi64>, window_strides = dense<1> : tensor<4xi64>} : (tensor<128x5x5x32xf32>, tensor<128x4x4x32xf32>, tensor<f32>) -> tensor<128x5x5x32xf32>  
+    }) {padding = dense<0> : tensor<4x2xi64>, window_dimensions = dense<[1, 2, 2, 1]> : tensor<4xi64>, window_strides = dense<1> : tensor<4xi64>} : (tensor<128x5x5x32xf32>, tensor<128x4x4x32xf32>, tensor<f32>) -> tensor<128x5x5x32xf32>
   return %0 : tensor<128x5x5x32xf32>
 }
\ No newline at end of file
diff --git a/libspu/compiler/tests/ops_negative.mlir b/libspu/compiler/tests/ops_negative.mlir
index f72fdbf8..2028c724 100644
--- a/libspu/compiler/tests/ops_negative.mlir
+++ b/libspu/compiler/tests/ops_negative.mlir
@@ -108,3 +108,18 @@ func.func @main(%arg0: tensor<10x!pphlo.pub<i32>>) -> (tensor<!pphlo.pub<i32>>)
   return %1 :  tensor<!pphlo.pub<i32>>
 }
 
+// -----
+
+func.func @main() -> tensor<!pphlo.pub<i32>> {
+    %0 = "pphlo.constant"() {value = dense<127> : tensor<i8>} : () -> tensor<!pphlo.pub<i8>>
+    %1 = "pphlo.slice"(%0) {limit_indices = dense<> : tensor<0xi64>, start_indices = dense<> : tensor<0xi64>, strides = dense<> : tensor<0xi64>} : (tensor<!pphlo.pub<i8>>) -> tensor<!pphlo.pub<i8>>
+    %2 = "pphlo.slice"(%1) {limit_indices = dense<> : tensor<0xi64>, start_indices = dense<> : tensor<0xi64>, strides = dense<> : tensor<0xi64>} : (tensor<!pphlo.pub<i8>>) -> tensor<!pphlo.pub<i8>>
+    %3 = "pphlo.slice"(%0) {limit_indices = dense<> : tensor<0xi64>, start_indices = dense<> : tensor<0xi64>, strides = dense<> : tensor<0xi64>} : (tensor<!pphlo.pub<i8>>) -> tensor<!pphlo.pub<i8>>
+    %4 = "pphlo.constant"() {value = dense<-1.7976931344453863E+308> : tensor<1x1xf64>} : () -> tensor<1x1x!pphlo.pub<f64>>
+    // expected-error @+1 {{op negative start index -9220555925398487041 in dimension 0}}
+    %5 = "pphlo.slice"(%4) {limit_indices = dense<[-9220555925398487041, 0]> : tensor<2xi64>, start_indices = dense<[-9220555925398487041, 0]> : tensor<2xi64>, strides = dense<[-9220555925398487041, 0]> : tensor<2xi64>} : (tensor<1x1x!pphlo.pub<f64>>) -> tensor<1x1x!pphlo.pub<f64>>
+    %6 = "pphlo.slice"(%4) {limit_indices = dense<[-8502447508339815911, -9223371558496411295]> : tensor<2xi64>, start_indices = dense<[-8502447508339815911, -9223371558496411295]> : tensor<2xi64>, strides = dense<[-8502447508339815911, -9223371558496411295]> : tensor<2xi64>} : (tensor<1x1x!pphlo.pub<f64>>) -> tensor<1x1x!pphlo.pub<f64>>
+    %7 = "pphlo.constant"() {value = dense<5> : tensor<i32>} : () -> tensor<!pphlo.pub<i32>>
+    "pphlo.return"(%7) : (tensor<!pphlo.pub<i32>>) -> ()
+  }
+
diff --git a/libspu/compiler/tests/optimize_denominator_with_bcst.mlir b/libspu/compiler/tests/optimize_denominator_with_bcst.mlir
new file mode 100644
index 00000000..47c9cd53
--- /dev/null
+++ b/libspu/compiler/tests/optimize_denominator_with_bcst.mlir
@@ -0,0 +1,11 @@
+// RUN: mlir-pphlo-opt --optimize-denominator-with-broadcast --split-input-file %s | FileCheck %s
+
+func.func @main(%arg0: tensor<16x!pphlo.sec<f32>>, %arg1: tensor<16x10000x!pphlo.sec<f32>>) -> (tensor<16x10000x!pphlo.sec<f32>>) {
+    //CHECK: %0 = "pphlo.reciprocal"(%arg0)
+    //CHECK: %1 = "pphlo.broadcast"(%0)
+    //CHECK: %2 = "pphlo.multiply"(%arg1, %1)
+    //CHECK: return %2
+    %0 = "pphlo.broadcast"(%arg0) {broadcast_dimensions = dense<0> : tensor<1xi64>} : (tensor<16x!pphlo.sec<f32>>) -> tensor<16x10000x!pphlo.sec<f32>>
+    %1 = "pphlo.divide"(%arg1, %0) : (tensor<16x10000x!pphlo.sec<f32>>, tensor<16x10000x!pphlo.sec<f32>>) -> tensor<16x10000x!pphlo.sec<f32>>
+    return %1 : tensor<16x10000x!pphlo.sec<f32>>
+}
diff --git a/libspu/compiler/tests/optimize_maxpool.mlir b/libspu/compiler/tests/optimize_maxpool.mlir
index 42153426..a8926d67 100644
--- a/libspu/compiler/tests/optimize_maxpool.mlir
+++ b/libspu/compiler/tests/optimize_maxpool.mlir
@@ -22,7 +22,7 @@ func.func @main(%arg0: tensor<129x24x24x16x!pphlo.sec<f32>>, %arg1: tensor<129x2
       %6 = "pphlo.add"(%arg2, %arg3) : (tensor<!pphlo.sec<f32>>, tensor<!pphlo.sec<f32>>) -> tensor<!pphlo.sec<f32>>
       "pphlo.return"(%6) : (tensor<!pphlo.sec<f32>>) -> ()
     }) {padding = dense<0> : tensor<4x2xi64>, window_dimensions = dense<[1, 2, 2, 1]> : tensor<4xi64>, window_strides = dense<1> : tensor<4xi64>} : (tensor<129x24x24x16x!pphlo.sec<f32>>, tensor<129x23x23x16x!pphlo.sec<f32>>, tensor<!pphlo.sec<f32>>) -> tensor<129x24x24x16x!pphlo.sec<f32>>
-   
+
     return %4, %5 : tensor<129x23x23x16x!pphlo.sec<f32>>, tensor<129x24x24x16x!pphlo.sec<f32>>
 }
 
@@ -43,6 +43,6 @@ func.func @main(%arg0: tensor<128x2x2x256x!pphlo.sec<f32>>, %arg1: tensor<128x1x
       %5 = "pphlo.add"(%arg2, %arg3) : (tensor<!pphlo.sec<f32>>, tensor<!pphlo.sec<f32>>) -> tensor<!pphlo.sec<f32>>
       "pphlo.return"(%5) : (tensor<!pphlo.sec<f32>>) -> ()
     }) {padding = dense<0> : tensor<4x2xi64>, window_dimensions = dense<[1, 2, 2, 1]> : tensor<4xi64>, window_strides = dense<1> : tensor<4xi64>} : (tensor<128x2x2x256x!pphlo.sec<f32>>, tensor<128x1x1x256x!pphlo.sec<f32>>, tensor<!pphlo.sec<f32>>) -> tensor<128x2x2x256x!pphlo.sec<f32>>
-   
+
     return %3, %4 : tensor<128x2x2x256x!pphlo.sec<f32>>, tensor<128x2x2x256x!pphlo.sec<f32>>
 }
diff --git a/libspu/compiler/tests/pphlo_type_inference_all_public.mlir b/libspu/compiler/tests/pphlo_type_inference_all_public.mlir
index 4b4b2bca..dcf4540e 100644
--- a/libspu/compiler/tests/pphlo_type_inference_all_public.mlir
+++ b/libspu/compiler/tests/pphlo_type_inference_all_public.mlir
@@ -1,7 +1,7 @@
 // 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>> 
+    // CHECK: "pphlo.constant"() {value = dense<1> : tensor<2x2xui8>} : () -> tensor<2x2x!pphlo.pub<ui8>>
     %0 = stablehlo.constant dense<1> : tensor<2x2xui8>
     // CHECK: %1 = "pphlo.sqrt"(%arg0) : (tensor<2x2x!pphlo.pub<f32>>) -> tensor<2x2x!pphlo.pub<f32>>
     %1 = "stablehlo.sqrt"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
diff --git a/libspu/compiler/tests/pphlo_type_inference_reduce.mlir b/libspu/compiler/tests/pphlo_type_inference_reduce.mlir
index d50af752..c687fa2c 100644
--- a/libspu/compiler/tests/pphlo_type_inference_reduce.mlir
+++ b/libspu/compiler/tests/pphlo_type_inference_reduce.mlir
@@ -1,8 +1,8 @@
 // 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>> { 
+// CHECK: func @main(%arg0: tensor<1024x1x!pphlo.sec<f32>>) -> tensor<1024x!pphlo.sec<f32>> {
 func.func @main(%arg1: tensor<1024x1xf32>) -> (tensor<1024xf32>) {
-    // CHECK: %0 = "pphlo.constant"() {value = dense<0.000000e+00> : tensor<f32>} : () -> tensor<!pphlo.pub<f32>> 
+    // CHECK: %0 = "pphlo.constant"() {value = dense<0.000000e+00> : tensor<f32>} : () -> tensor<!pphlo.pub<f32>>
     %0 = "stablehlo.constant"() {value = dense<0.000000e+00> : tensor<f32>} : () -> tensor<f32>
     // CHECK: %2 = "pphlo.reduce"(%arg0, %1) ({
     // CHECK:       ^bb0(%arg1: tensor<!pphlo.sec<f32>>, %arg2: tensor<!pphlo.sec<f32>>):
diff --git a/libspu/core/encoding_test.cc b/libspu/core/encoding_test.cc
index 8df2aeda..576a18dd 100644
--- a/libspu/core/encoding_test.cc
+++ b/libspu/core/encoding_test.cc
@@ -32,8 +32,8 @@ TEST(EncodingTypeTest, EncodeDecodeMap) {
   EXPECT_EQ(getEncodeType(PT_I64), DT_I64);
   EXPECT_EQ(getEncodeType(PT_U64), DT_U64);
   EXPECT_EQ(getEncodeType(PT_BOOL), DT_I1);
-  EXPECT_EQ(getEncodeType(PT_F32), DT_FXP);
-  EXPECT_EQ(getEncodeType(PT_F64), DT_FXP);
+  EXPECT_EQ(getEncodeType(PT_F32), DT_F32);
+  EXPECT_EQ(getEncodeType(PT_F64), DT_F64);
 
   EXPECT_EQ(getDecodeType(DT_I8), PT_I8);
   EXPECT_EQ(getDecodeType(DT_U8), PT_U8);
@@ -44,7 +44,8 @@ TEST(EncodingTypeTest, EncodeDecodeMap) {
   EXPECT_EQ(getDecodeType(DT_I64), PT_I64);
   EXPECT_EQ(getDecodeType(DT_U64), PT_U64);
   EXPECT_EQ(getDecodeType(DT_I1), PT_BOOL);
-  EXPECT_EQ(getDecodeType(DT_FXP), PT_F32);
+  EXPECT_EQ(getDecodeType(DT_F32), PT_F32);
+  EXPECT_EQ(getDecodeType(DT_F64), PT_F64);
 }
 
 using Field64 = std::integral_constant<FieldType, FM64>;
@@ -107,17 +108,25 @@ TYPED_TEST(FloatEncodingTest, Works) {
 
   DataType encoded_dtype;
   auto encoded = encodeToRing(frm, kField, kFxpBits, &encoded_dtype);
-  EXPECT_EQ(encoded_dtype, DT_FXP);
 
+  if constexpr (std::is_same_v<FloatT, float>) {
+    EXPECT_EQ(encoded_dtype, DT_F32);
+  } else {
+    EXPECT_EQ(encoded_dtype, DT_F64);
+  }
   PtType out_pt_type;
   auto decoded = decodeFromRing(encoded, encoded_dtype, kFxpBits, &out_pt_type);
-  EXPECT_EQ(out_pt_type, PT_F32);
 
-  float* out_ptr = &decoded.at<float>(0);
+  if constexpr (std::is_same_v<FloatT, float>) {
+    EXPECT_EQ(encoded_dtype, DT_F32);
+  } else {
+    EXPECT_EQ(encoded_dtype, DT_F64);
+  }
+  auto* out_ptr = &decoded.at<FloatT>(0);
   const int64_t kReprBits = SizeOf(kField) * 8 - 2;
   const int64_t kScale = 1LL << kFxpBits;
-  EXPECT_EQ(out_ptr[0], -static_cast<float>((1LL << kReprBits)) / kScale);
-  EXPECT_EQ(out_ptr[1], static_cast<float>((1LL << kReprBits) - 1) / kScale);
+  EXPECT_EQ(out_ptr[0], -static_cast<FloatT>((1LL << kReprBits)) / kScale);
+  EXPECT_EQ(out_ptr[1], static_cast<FloatT>((1LL << kReprBits) - 1) / kScale);
   EXPECT_EQ(out_ptr[2], -1.0);
   EXPECT_EQ(out_ptr[3], 0.0);
   EXPECT_EQ(out_ptr[4], 1.0);
diff --git a/libspu/core/value_test.cc b/libspu/core/value_test.cc
index 57169d57..3e65fdaa 100644
--- a/libspu/core/value_test.cc
+++ b/libspu/core/value_test.cc
@@ -50,7 +50,7 @@ TEST(ValueTest, Empty) {
 //    EXPECT_EQ(a.GetBuffer()->data<int>()[0], 1);
 //  }
 
-//  Value b(VIS_SECRET, DT_FXP, {}, yacl::Buffer(&dval, sizeof(double)));
+//  Value b(VIS_SECRET, DT_F32, {}, yacl::Buffer(&dval, sizeof(double)));
 //  {
 //    b.GetBuffer()->data<double>()[0] = 1;
 
@@ -58,7 +58,7 @@ TEST(ValueTest, Empty) {
 //    EXPECT_FALSE(b.isPublic());
 //    EXPECT_TRUE(b.isSecret());
 
-//    EXPECT_EQ(b.dtype(), DT_FXP);
+//    EXPECT_EQ(b.dtype(), DT_F32);
 //    EXPECT_TRUE(b.isFxp());
 //    EXPECT_FALSE(b.isInt());
 //    EXPECT_EQ(b.GetBuffer()->data<double>()[0], 1);
@@ -84,7 +84,7 @@ TEST(ValueTest, Empty) {
 //    EXPECT_FALSE(d.isPublic());
 //    EXPECT_TRUE(d.isSecret());
 
-//    EXPECT_EQ(d.dtype(), DT_FXP);
+//    EXPECT_EQ(d.dtype(), DT_F32);
 //    EXPECT_TRUE(d.isFxp());
 //    EXPECT_FALSE(d.isInt());
 
diff --git a/libspu/device/pphlo/pphlo_executor_test.cc b/libspu/device/pphlo/pphlo_executor_test.cc
index 3206c01c..2299c5c2 100644
--- a/libspu/device/pphlo/pphlo_executor_test.cc
+++ b/libspu/device/pphlo/pphlo_executor_test.cc
@@ -25,102 +25,11 @@
 #include "gtest/gtest.h"
 #include "xtensor/xrandom.hpp"
 
-#include "libspu/compiler/common/compilation_context.h"
-#include "libspu/compiler/compile.h"
-#include "libspu/device/api.h"
+#include "libspu/device/pphlo/pphlo_executor_test_runner.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"
-
-namespace spu::device {
-namespace {
-
-class Runner {
- public:
-  Runner(size_t world_size, FieldType field, ProtocolKind protocol)
-      : world_size_(world_size) {
-    config_.set_field(field);
-    config_.set_protocol(protocol);
-    config_.set_enable_type_checker(true);
-    io_ = std::make_unique<LocalIo>(world_size_, config_);
-  }
-
-  auto &getConfig() { return config_; }
-
-  template <typename T>
-  void addInput(const T &input, Visibility vis = Visibility::VIS_PUBLIC) {
-    const std::string name = fmt::format("input{}", input_idx_++);
-    io_->InFeed(name, input, vis);
-    executable_.add_input_names(name);
-  }
-
-  static std::string compileMHlo(const std::string &mhlo,
-                                 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);
-    }
-
-    compiler::CompilationContext ctx;
-    return compiler::compile(&ctx, source.SerializeAsString());
-  }
-
-  void run(const std::string &mlir, size_t num_output = 1) {
-    for (size_t idx = 0; idx < num_output; ++idx) {
-      executable_.add_output_names(fmt::format("output{}", idx));
-    }
-    executable_.set_code(mlir);
-    ::spu::mpc::utils::simulate(
-        world_size_, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
-          RuntimeConfig conf;
-          conf.CopyFrom(config_);
-          if (lctx->Rank() == 0) {
-            // conf.set_enable_action_trace(true);
-          }
-          SPUContext sctx = kernel::test::makeSPUContext(conf, lctx);
-          auto *env = io_->GetSymbolTable(lctx->Rank());
-          pphlo::PPHloExecutor executor;
-          execute(&executor, &sctx, executable_, env);
-        });
-  }
-
-  template <typename T>
-  void verifyOutput(const T *expected, size_t idx = 0) {
-    const auto &out = io_->OutFeed(fmt::format("output{}", idx));
-
-    size_t numel = out.numel();
-    const auto *in_ptr = static_cast<const T *>(out.data());
-
-    // TODO: handle strides
-    for (size_t i = 0; i < numel; ++i) {
-      if constexpr (std::is_integral_v<T>) {
-        EXPECT_EQ(in_ptr[i], expected[i]) << "i = " << i << "\n";
-      } else {
-        EXPECT_TRUE(std::abs(in_ptr[i] - expected[i]) <= 1e-2)
-            << "i = " << i << " in = " << in_ptr[i]
-            << " expected = " << expected[i] << "\n";
-      }
-    }
-  }
-
-  template <typename T, std::enable_if_t<std::is_scalar_v<T>, bool> = true>
-  void verifyScalarOutput(T expected, size_t idx = 0) {
-    verifyOutput(&expected, idx);
-  }
-
- private:
-  size_t world_size_;
-  RuntimeConfig config_;
-  std::unique_ptr<LocalIo> io_;
-  size_t input_idx_{0};
-  ExecutableProto executable_;
-};
 
-}  // namespace
+namespace spu::device::pphlo::test {
 
 class ExecutorTest : public ::testing::TestWithParam<
                          std::tuple<size_t, FieldType, ProtocolKind>> {};
@@ -938,7 +847,7 @@ TEST_P(ExecutorTest, Simple4x4Conv2DWith2x2Kernel) {
   }}};
   r.addInput(rhs);
 
-  auto ir = spu::device::Runner::compileMHlo(R"(
+  auto ir = r.compileMHlo(R"(
 func.func @main(%arg0: tensor<1x1x4x4xf32>, %arg1: tensor<1x1x2x2xf32>) -> (tensor<1x1x4x4xf32>) {
     %0 = mhlo.convolution(%arg0, %arg1)
             dim_numbers = [b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1],
@@ -949,7 +858,7 @@ func.func @main(%arg0: tensor<1x1x4x4xf32>, %arg1: tensor<1x1x2x2xf32>) -> (tens
             } : (tensor<1x1x4x4xf32>, tensor<1x1x2x2xf32>) -> tensor<1x1x4x4xf32>
     return %0 : tensor<1x1x4x4xf32>
 })",
-                                             {VIS_PUBLIC, VIS_PUBLIC});
+                          {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -977,7 +886,7 @@ TEST_P(ExecutorTest, Conv2DGeneralDimensions) {
 
   r.addInput(rhs);
 
-  auto ir = spu::device::Runner::compileMHlo(R"(
+  auto ir = r.compileMHlo(R"(
 func.func @main(%arg0: tensor<2x3x1x4xf32>, %arg1:tensor<1x3x2x3xf32>) -> (tensor<1x1x1x2xf32>) {
     %0 = mhlo.convolution(%arg0, %arg1)
           dim_numbers = [f, 0, b, 1]x[o, 1, i,0]->[f, 0, b, 1],
@@ -988,7 +897,7 @@ func.func @main(%arg0: tensor<2x3x1x4xf32>, %arg1:tensor<1x3x2x3xf32>) -> (tenso
           } : (tensor<2x3x1x4xf32>,tensor<1x3x2x3xf32>) -> tensor<1x1x1x2xf32>
     return %0 : tensor<1x1x1x2xf32>
 })",
-                                             {VIS_PUBLIC, VIS_PUBLIC});
+                          {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -1015,7 +924,7 @@ TEST_P(ExecutorTest, DilatedBaseConv2DWithHighPadding) {
 
   r.addInput(rhs);
 
-  auto ir = spu::device::Runner::compileMHlo(R"(
+  auto ir = r.compileMHlo(R"(
 func.func @main(%arg0: tensor<1x1x4x4xf32>, %arg1: tensor<1x1x2x2xf32>) -> (tensor<1x1x7x7xf32>) {
     %0 = mhlo.convolution(%arg0, %arg1)
           dim_numbers = [b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1],
@@ -1026,7 +935,7 @@ func.func @main(%arg0: tensor<1x1x4x4xf32>, %arg1: tensor<1x1x2x2xf32>) -> (tens
           } : (tensor<1x1x4x4xf32>, tensor<1x1x2x2xf32>) -> tensor<1x1x7x7xf32>
     return %0 : tensor<1x1x7x7xf32>
 })",
-                                             {VIS_PUBLIC, VIS_PUBLIC});
+                          {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -1059,7 +968,7 @@ TEST_P(ExecutorTest, DilatedBaseConv2DWithLowAndHighPadding) {
 
   r.addInput(rhs);
 
-  auto ir = spu::device::Runner::compileMHlo(R"(
+  auto ir = r.compileMHlo(R"(
 func.func @main(%arg0: tensor<1x1x4x4xf32>, %arg1: tensor<1x1x2x2xf32>) -> (tensor<1x1x8x8xf32>) {
     %0 = mhlo.convolution(%arg0, %arg1)
           dim_numbers = [b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1],
@@ -1070,7 +979,7 @@ func.func @main(%arg0: tensor<1x1x4x4xf32>, %arg1: tensor<1x1x2x2xf32>) -> (tens
           } : (tensor<1x1x4x4xf32>, tensor<1x1x2x2xf32>) -> tensor<1x1x8x8xf32>
     return %0 : tensor<1x1x8x8xf32>
 })",
-                                             {VIS_PUBLIC, VIS_PUBLIC});
+                          {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -1105,7 +1014,7 @@ TEST_P(ExecutorTest, FlatRhsDilation) {
 
   r.addInput(rhs);
 
-  auto ir = spu::device::Runner::compileMHlo(R"(
+  auto ir = r.compileMHlo(R"(
 func.func @main(%arg0: tensor<1x1x4x6xf32>, %arg1: tensor<1x1x2x3xf32>) -> (tensor<1x1x2x2xf32>) {
     %0 = mhlo.convolution(%arg0, %arg1)
           dim_numbers = [b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1],
@@ -1116,7 +1025,7 @@ func.func @main(%arg0: tensor<1x1x4x6xf32>, %arg1: tensor<1x1x2x3xf32>) -> (tens
           } : (tensor<1x1x4x6xf32>, tensor<1x1x2x3xf32>) -> tensor<1x1x2x2xf32>
     return %0 : tensor<1x1x2x2xf32>
 })",
-                                             {VIS_PUBLIC, VIS_PUBLIC});
+                          {VIS_PUBLIC, VIS_PUBLIC});
 
   r.run(ir);
 
@@ -1954,7 +1863,7 @@ TEST_P(ExecutorTest, OptimizedMaxPool1) {
       {3, 1}   //
   });
 
-  auto ir = spu::device::Runner::compileMHlo(R"(
+  auto ir = r.compileMHlo(R"(
 func.func @main(%arg0: tensor<4x6xi32>, %arg1: tensor<2x2xi32>) -> (tensor<2x2xi32>, tensor<4x6xi32>) {
   %0 = mhlo.constant dense<0> : tensor<i32>
   %1 = "mhlo.reduce_window"(%arg0, %0) ({
@@ -1974,7 +1883,7 @@ func.func @main(%arg0: tensor<4x6xi32>, %arg1: tensor<2x2xi32>) -> (tensor<2x2xi
     }) {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});
+                          {VIS_PUBLIC, VIS_PUBLIC});
 
   EXPECT_THAT(ir, testing::HasSubstr("pphlo.maxpool_scatter"));
 
@@ -2202,7 +2111,7 @@ TEST_P(ExecutorTest, MixedPayload) {
 
   r.addInput(op, VIS_SECRET);
 
-  r.run(spu::device::Runner::compileMHlo(
+  r.run(r.compileMHlo(
             R"(
 func.func @main(%arg0: tensor<20xi32>) -> (tensor<20xi32>, tensor<20xi32>) {
     %0 = "mhlo.iota"() {iota_dimension = 0 : i64} : () -> tensor<20xi32>
@@ -2241,4 +2150,4 @@ INSTANTIATE_TEST_SUITE_P(
                          std::get<2>(p.param));
     });
 
-}  // namespace spu::device
+}  // namespace spu::device::pphlo::test
diff --git a/libspu/device/pphlo/pphlo_executor_test_runner.cc b/libspu/device/pphlo/pphlo_executor_test_runner.cc
new file mode 100644
index 00000000..bad6ec3c
--- /dev/null
+++ b/libspu/device/pphlo/pphlo_executor_test_runner.cc
@@ -0,0 +1,66 @@
+// 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/device/pphlo/pphlo_executor_test_runner.h"
+
+#include "libspu/compiler/common/compilation_context.h"
+#include "libspu/compiler/compile.h"
+#include "libspu/device/api.h"
+#include "libspu/device/pphlo/pphlo_executor.h"
+#include "libspu/kernel/test_util.h"
+#include "libspu/mpc/utils/simulate.h"
+
+namespace spu::device::pphlo::test {
+
+Runner::Runner(size_t world_size, FieldType field, ProtocolKind protocol)
+    : world_size_(world_size) {
+  config_.set_field(field);
+  config_.set_protocol(protocol);
+  config_.set_enable_type_checker(true);
+  io_ = std::make_unique<LocalIo>(world_size_, config_);
+}
+
+std::string Runner::compileMHlo(const std::string &mhlo,
+                                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);
+  }
+
+  compiler::CompilationContext ctx;
+  return compiler::compile(&ctx, source.SerializeAsString());
+}
+
+void Runner::run(const std::string &mlir, size_t num_output) {
+  for (size_t idx = 0; idx < num_output; ++idx) {
+    executable_.add_output_names(fmt::format("output{}", idx));
+  }
+  executable_.set_code(mlir);
+  ::spu::mpc::utils::simulate(
+      world_size_, [&](const std::shared_ptr<yacl::link::Context> &lctx) {
+        RuntimeConfig conf;
+        conf.CopyFrom(config_);
+        if (lctx->Rank() == 0) {
+          // conf.set_enable_action_trace(true);
+        }
+        SPUContext sctx = kernel::test::makeSPUContext(conf, lctx);
+        auto *env = io_->GetSymbolTable(lctx->Rank());
+        pphlo::PPHloExecutor executor;
+        execute(&executor, &sctx, executable_, env);
+      });
+}
+
+}  // namespace spu::device::pphlo::test
diff --git a/libspu/device/pphlo/pphlo_executor_test_runner.h b/libspu/device/pphlo/pphlo_executor_test_runner.h
new file mode 100644
index 00000000..3b52cbe5
--- /dev/null
+++ b/libspu/device/pphlo/pphlo_executor_test_runner.h
@@ -0,0 +1,75 @@
+// 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 "fmt/format.h"
+
+#include "libspu/device/test_utils.h"
+
+#include "libspu/spu.pb.h"
+
+namespace spu::device::pphlo::test {
+
+class Runner {
+ public:
+  Runner(size_t world_size, FieldType field, ProtocolKind protocol);
+
+  auto &getConfig() { return config_; }
+
+  template <typename T>
+  void addInput(const T &input, Visibility vis = Visibility::VIS_PUBLIC) {
+    const std::string name = fmt::format("input{}", input_idx_++);
+    io_->InFeed(name, input, vis);
+    executable_.add_input_names(name);
+  }
+
+  std::string compileMHlo(const std::string &mhlo,
+                          const std::vector<spu::Visibility> &vis);
+
+  void run(const std::string &mlir, size_t num_output = 1);
+
+  template <typename T>
+  void verifyOutput(const T *expected, size_t idx = 0) {
+    const auto &out = io_->OutFeed(fmt::format("output{}", idx));
+
+    size_t numel = out.numel();
+    const auto *in_ptr = static_cast<const T *>(out.data());
+
+    // TODO: handle strides
+    for (size_t i = 0; i < numel; ++i) {
+      if constexpr (std::is_integral_v<T>) {
+        EXPECT_EQ(in_ptr[i], expected[i]) << "i = " << i << "\n";
+      } else {
+        EXPECT_TRUE(std::abs(in_ptr[i] - expected[i]) <= 1e-2)
+            << "i = " << i << " in = " << in_ptr[i]
+            << " expected = " << expected[i] << "\n";
+      }
+    }
+  }
+
+  template <typename T, std::enable_if_t<std::is_scalar_v<T>, bool> = true>
+  void verifyScalarOutput(T expected, size_t idx = 0) {
+    verifyOutput(&expected, idx);
+  }
+
+ private:
+  size_t world_size_;
+  RuntimeConfig config_;
+  std::unique_ptr<LocalIo> io_;
+  size_t input_idx_{0};
+  ExecutableProto executable_;
+};
+
+}  // namespace spu::device::pphlo::test
diff --git a/libspu/kernel/hal/constants_test.cc b/libspu/kernel/hal/constants_test.cc
index 7fd3ef67..238fe8d0 100644
--- a/libspu/kernel/hal/constants_test.cc
+++ b/libspu/kernel/hal/constants_test.cc
@@ -31,7 +31,7 @@ TEST(ConstantsTest, Scalar) {
   EXPECT_TRUE(i.isPublic());
   EXPECT_TRUE(i.isInt());
 
-  Value f = constant(&ctx, 0.0, DT_FXP);
+  Value f = constant(&ctx, 0.0F, DT_F32);
   EXPECT_TRUE(f.shape().empty());
   EXPECT_TRUE(f.strides().empty());
   EXPECT_EQ(f.numel(), 1);
@@ -43,7 +43,7 @@ TEST(ConstantsTest, Tensor) {
   SPUContext ctx = test::makeSPUContext();
 
   xt::xarray<float> raw = {1.0F};
-  Value x = constant(&ctx, raw, DT_FXP);
+  Value x = constant(&ctx, raw, DT_F32);
   EXPECT_THAT(x.shape(), testing::ElementsAre(1));
   EXPECT_THAT(x.strides(), testing::ElementsAre(0));
   EXPECT_EQ(x.numel(), 1);
@@ -58,7 +58,7 @@ TEST(ConstantsTest, TensorBroadcast) {
       {1.0, 2.0},
   };
 
-  Value x = constant(&ctx, raw, DT_FXP, {6, 2});
+  Value x = constant(&ctx, raw, DT_F32, {6, 2});
   EXPECT_THAT(x.shape(), testing::ElementsAre(6, 2));
   EXPECT_THAT(x.strides(), testing::ElementsAre(0, 1));
   EXPECT_EQ(x.numel(), 12);
@@ -79,8 +79,8 @@ TEST(ConstantsTest, Initializer) {
   // EXPECT_EQ(constant(&ctx, 0, DT_U32).dtype(), DT_U32); // FIXME
   // EXPECT_EQ(constant(&ctx, 0, DT_I64).dtype(), DT_I64); // FIXME
   // EXPECT_EQ(constant(&ctx, 0, DT_U64).dtype(), DT_U64); // FIXME
-  // EXPECT_EQ(constant(&ctx, 0, DT_FXP).dtype(), DT_FXP); // FIXME
-  EXPECT_EQ(constant(&ctx, 0.0, DT_FXP).dtype(), DT_FXP);
+  EXPECT_EQ(constant(&ctx, 0.0F, DT_F32).dtype(), DT_F32);
+  EXPECT_EQ(constant(&ctx, 0.0, DT_F64).dtype(), DT_F64);
 }
 
 }  // namespace spu::kernel::hal
diff --git a/libspu/kernel/hal/fxp_approx_test.cc b/libspu/kernel/hal/fxp_approx_test.cc
index 64b5f0ae..f8b9da39 100644
--- a/libspu/kernel/hal/fxp_approx_test.cc
+++ b/libspu/kernel/hal/fxp_approx_test.cc
@@ -18,6 +18,7 @@
 #include "xtensor/xio.hpp"
 
 #include "libspu/kernel/hal/constants.h"
+#include "libspu/kernel/hal/type_cast.h"
 #include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
@@ -31,9 +32,9 @@ TEST(FxpTest, ExponentialPublic) {
       6.7,    8.0,    10.5,  12.5,  14.3,  16.7,  18.0, 20.0,
   };
 
-  Value a = constant(&ctx, x, DT_FXP);
+  Value a = constant(&ctx, x, DT_F32);
   Value c = f_exp(&ctx, a);
-  EXPECT_EQ(c.dtype(), DT_FXP);
+  EXPECT_EQ(c.dtype(), DT_F32);
 
   auto y = dump_public_as<float>(&ctx, c);
   EXPECT_TRUE(xt::allclose(xt::exp(x), y, 0.01, 0.001))
@@ -54,9 +55,9 @@ TEST(FxpTest, ExponentialTaylorSeries) {
 
   Value a = test::makeValue(&ctx, x, VIS_SECRET);
   Value c = detail::exp_taylor_series(&ctx, a);
-  EXPECT_EQ(c.dtype(), DT_FXP);
+  EXPECT_EQ(c.dtype(), DT_F32);
 
-  auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+  auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
   EXPECT_TRUE(xt::allclose(xt::exp(x), y, 0.01, 0.001))
       << xt::exp(x) << std::endl
       << y;
@@ -69,9 +70,9 @@ TEST(FxpTest, ExponentialPade) {
 
   Value a = test::makeValue(&ctx, x, VIS_SECRET);
   Value c = detail::exp_pade_approx(&ctx, a);
-  EXPECT_EQ(c.dtype(), DT_FXP);
+  EXPECT_EQ(c.dtype(), DT_F32);
 
-  auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+  auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
   EXPECT_TRUE(xt::allclose(xt::exp(x), y, 0.01, 0.001))
       << xt::exp(x) << std::endl
       << y;
@@ -84,9 +85,9 @@ TEST(FxpTest, Log) {
   xt::xarray<float> x = {{0.05, 0.5}, {5, 50}};
   // public log
   {
-    Value a = constant(&ctx, x, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
     Value c = f_log(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto y = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(xt::log(x), y, 0.01, 0.001))
@@ -98,9 +99,9 @@ TEST(FxpTest, Log) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_log(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     // low precision
     EXPECT_TRUE(xt::allclose(xt::log(x), y, 0.01, 0.001))
         << xt::log(x) << std::endl
@@ -115,9 +116,9 @@ TEST(FxpTest, Log2) {
   xt::xarray<float> x = {{0.05, 0.5}, {5, 50}};
   // public log
   {
-    Value a = constant(&ctx, x, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
     Value c = f_log2(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto y = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(xt::log2(x), y, 0.01, 0.001))
@@ -129,9 +130,9 @@ TEST(FxpTest, Log2) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_log2(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     // low precision
     EXPECT_TRUE(xt::allclose(xt::log2(x), y, 0.01, 0.001))
         << xt::log2(x) << std::endl
@@ -147,9 +148,9 @@ TEST(FxpTest, Log1p) {
 
   // public log1p
   {
-    Value a = constant(&ctx, x, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
     Value c = f_log1p(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto y = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(xt::log1p(x), y, 0.01, 0.001))
@@ -160,9 +161,9 @@ TEST(FxpTest, Log1p) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_log1p(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     // low precision
     EXPECT_TRUE(xt::allclose(xt::log1p(x), y, 0.01, 0.001))
         << xt::log1p(x) << std::endl
@@ -184,9 +185,9 @@ TEST(FxpTest, Exp2) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_exp2(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     EXPECT_TRUE(xt::allclose(xt::exp2(x), y, 0.01, 0.001))
         << xt::exp2(x) << std::endl
         << y;
@@ -204,9 +205,9 @@ TEST(FxpTest, Tanh) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_tanh(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     EXPECT_TRUE(xt::allclose(xt::tanh(x), y, 0.01, 0.001))
         << xt::tanh(x) << std::endl
         << y;
@@ -224,9 +225,9 @@ TEST(FxpTest, Rsqrt) {
 
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_rsqrt(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     EXPECT_TRUE(xt::allclose(expected_y, y, 0.01, 0.001))
         << expected_y << std::endl
         << y;
@@ -242,9 +243,9 @@ TEST(FxpTest, Rsqrt) {
 
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_rsqrt(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     EXPECT_TRUE(xt::allclose(expected_y, y, 0.01, 0.001))
         << expected_y << std::endl
         << y;
@@ -261,9 +262,9 @@ TEST(FxpTest, Rsqrt) {
     xt::xarray<float> x = xt::random::rand<float>({200, 1}, 256, 4096);
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_rsqrt(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto z = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto z = dump_public_as<float>(&ctx, reveal(&ctx, c));
     auto e = 1.0F / xt::sqrt(x);
     auto r = xt::abs(z - e) / e * 100;
     auto mm = xt::minmax(r)();
@@ -282,9 +283,9 @@ TEST(FxpTest, Sqrt) {
 
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_sqrt(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     EXPECT_TRUE(xt::allclose(expected_y, y, 0.01, 0.001))
         << expected_y << std::endl
         << y;
@@ -300,9 +301,9 @@ TEST(FxpTest, Sqrt) {
 
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_sqrt(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     EXPECT_TRUE(xt::allclose(expected_y, y, 0.01, 0.001))
         << expected_y << std::endl
         << y;
diff --git a/libspu/kernel/hal/fxp_base_test.cc b/libspu/kernel/hal/fxp_base_test.cc
index 8ed1dad5..9eed4132 100644
--- a/libspu/kernel/hal/fxp_base_test.cc
+++ b/libspu/kernel/hal/fxp_base_test.cc
@@ -19,6 +19,7 @@
 
 #include "libspu/core/parallel_utils.h"
 #include "libspu/kernel/hal/constants.h"
+#include "libspu/kernel/hal/type_cast.h"
 #include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hal {
@@ -80,10 +81,10 @@ TEST_P(FxpMmulTest, Works) {
   auto t_z = xarrayMMul(x, y);
 
   {  // public
-    Value a = constant(&ctx, x, DT_FXP);
-    Value b = constant(&ctx, y, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
+    Value b = constant(&ctx, y, DT_F32);
     Value c = f_mmul(&ctx, a, b);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto z = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(t_z, z, 0.01, 0.001)) << t_z << std::endl << z;
@@ -93,9 +94,9 @@ TEST_P(FxpMmulTest, Works) {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value b = test::makeValue(&ctx, y, VIS_SECRET);
     Value c = f_mmul(&ctx, a, b);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto z = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto z = dump_public_as<float>(&ctx, reveal(&ctx, c));
     EXPECT_TRUE(xt::allclose(t_z, z, 0.01, 0.001)) << t_z << std::endl << z;
   }
 }
@@ -113,9 +114,9 @@ TEST(FxpTest, Reciprocal) {
 
   // public reciprocal
   {
-    Value a = constant(&ctx, x, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
     Value c = f_reciprocal(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto y = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(1.0F / x, y, 0.001, 0.0001))
@@ -127,9 +128,9 @@ TEST(FxpTest, Reciprocal) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_reciprocal(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     EXPECT_TRUE(xt::allclose(1.0F / x, y, 0.001, 0.0001))
         << (1.0 / x) << std::endl
         << y;
@@ -145,10 +146,10 @@ TEST(FxpTest, Div) {
 
   // public div
   {
-    Value a = constant(&ctx, x, DT_FXP);
-    Value b = constant(&ctx, y, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
+    Value b = constant(&ctx, y, DT_F32);
     Value c = f_div(&ctx, a, b);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto z = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(x / y, z, 0.001, 0.0001)) << (x / y) << std::endl
@@ -160,9 +161,9 @@ TEST(FxpTest, Div) {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value b = test::makeValue(&ctx, y, VIS_SECRET);
     Value c = f_div(&ctx, a, b);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto z = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto z = dump_public_as<float>(&ctx, reveal(&ctx, c));
     EXPECT_TRUE(xt::allclose(x / y, z, 0.01, 0.001)) << (x / y) << std::endl
                                                      << z;
   }
@@ -173,9 +174,9 @@ TEST(FxpTest, Div) {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value b = test::makeValue(&ctx, y, VIS_SECRET);
     Value c = f_div(&ctx, a, b);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto z = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto z = dump_public_as<float>(&ctx, reveal(&ctx, c));
     auto e = x / y;
     auto r = xt::abs(z - e) / e * 100;
     auto mm = xt::minmax(r)();
@@ -191,9 +192,9 @@ TEST(FxpTest, Abs) {
 
   // public abs
   {
-    Value a = constant(&ctx, x, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
     Value c = f_abs(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto y = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(xt::abs(x), y, 0.01, 0.05))
@@ -205,9 +206,9 @@ TEST(FxpTest, Abs) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_abs(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     // low precision
     EXPECT_TRUE(xt::allclose(xt::abs(x), y, 0.1, 0.5))
         << xt::abs(x) << std::endl
@@ -223,9 +224,9 @@ TEST(FxpTest, Floor) {
 
   // public floor
   {
-    Value a = constant(&ctx, x, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
     Value c = f_floor(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto y = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(xt::floor(x), y, 0.01, 0.001))
@@ -237,9 +238,9 @@ TEST(FxpTest, Floor) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_floor(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     // low precision
     EXPECT_TRUE(xt::allclose(xt::floor(x), y, 0.01, 0.001))
         << xt::floor(x) << std::endl
@@ -255,9 +256,9 @@ TEST(FxpTest, Ceil) {
 
   // public ceil
   {
-    Value a = constant(&ctx, x, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
     Value c = f_ceil(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto y = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(xt::ceil(x), y, 0.01, 0.001))
@@ -269,9 +270,9 @@ TEST(FxpTest, Ceil) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = f_ceil(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     // low precision
     EXPECT_TRUE(xt::allclose(xt::ceil(x), y, 0.01, 0.001))
         << xt::ceil(x) << std::endl
diff --git a/libspu/kernel/hal/polymorphic_test.cc b/libspu/kernel/hal/polymorphic_test.cc
index 7119d60e..f2797d45 100644
--- a/libspu/kernel/hal/polymorphic_test.cc
+++ b/libspu/kernel/hal/polymorphic_test.cc
@@ -24,6 +24,7 @@
 #include "xtensor/xvectorize.hpp"
 
 #include "libspu/kernel/hal/constants.h"
+#include "libspu/kernel/hal/type_cast.h"
 #include "libspu/kernel/test_util.h"
 #include "libspu/mpc/utils/linalg.h"
 
@@ -440,6 +441,17 @@ template <typename S>
 class MathUnaryTest : public ::testing::Test {};
 TYPED_TEST_SUITE(MathUnaryTest, MathUnaryTestTypes);
 
+using FpOnlyMathUnaryTestTypes = ::testing::Types<
+    // s
+    std::tuple<float, secret_v, float>,  // (sfxp)
+    // p
+    std::tuple<float, public_v, float>  // (pfxp)
+    >;
+
+template <typename S>
+class FpOnlyMathUnaryTest : public ::testing::Test {};
+TYPED_TEST_SUITE(FpOnlyMathUnaryTest, FpOnlyMathUnaryTestTypes);
+
 TYPED_TEST(MathUnaryTest, Negate) {
   using IN_DT = typename std::tuple_element<0, TypeParam>::type;
   using IN_VT = typename std::tuple_element<1, TypeParam>::type;
@@ -470,7 +482,7 @@ TYPED_TEST(MathUnaryTest, Abs) {
   EXPECT_TRUE(xt::allclose(xt::abs(x), z, 0.01, 0.001));
 }
 
-TYPED_TEST(MathUnaryTest, Exp) {
+TYPED_TEST(FpOnlyMathUnaryTest, Exp) {
   using IN_DT = typename std::tuple_element<0, TypeParam>::type;
   using IN_VT = typename std::tuple_element<1, TypeParam>::type;
   using RES_DT = float;
@@ -528,7 +540,7 @@ TYPED_TEST(MathUnaryTest, Reciprocal) {
   EXPECT_TRUE(xt::allclose(log_x, z, 0.1, 0.05)) << log_x << std::endl << z;
 }
 
-TYPED_TEST(MathUnaryTest, Log) {
+TYPED_TEST(FpOnlyMathUnaryTest, Log) {
   using IN_DT = typename std::tuple_element<0, TypeParam>::type;
   using IN_VT = typename std::tuple_element<1, TypeParam>::type;
   using RES_DT = float;
@@ -568,7 +580,7 @@ TYPED_TEST(MathUnaryTest, Logistic) {
       << z;
 }
 
-TYPED_TEST(MathUnaryTest, Tanh) {
+TYPED_TEST(FpOnlyMathUnaryTest, Tanh) {
   using IN_DT = typename std::tuple_element<0, TypeParam>::type;
   using IN_VT = typename std::tuple_element<1, TypeParam>::type;
   using RES_DT = float;
@@ -585,7 +597,7 @@ TYPED_TEST(MathUnaryTest, Tanh) {
       << z;
 }
 
-TYPED_TEST(MathUnaryTest, SqrtInv) {
+TYPED_TEST(FpOnlyMathUnaryTest, RSqrt) {
   using IN_DT = typename std::tuple_element<0, TypeParam>::type;
   using IN_VT = typename std::tuple_element<1, TypeParam>::type;
   using RES_DT = float;
@@ -603,7 +615,7 @@ TYPED_TEST(MathUnaryTest, SqrtInv) {
       << y;
 }
 
-TYPED_TEST(MathUnaryTest, Sqrt) {
+TYPED_TEST(FpOnlyMathUnaryTest, Sqrt) {
   using IN_DT = typename std::tuple_element<0, TypeParam>::type;
   using IN_VT = typename std::tuple_element<1, TypeParam>::type;
   using RES_DT = float;
@@ -656,9 +668,9 @@ TEST_P(LogisticTest, Logistic) {
 
   // public logistic
   {
-    Value a = constant(&ctx, x, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
     Value c = logistic(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto y = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(1.0 / (1.0 + xt::exp(-x)), y, 0.1, 0.5))
@@ -670,9 +682,9 @@ TEST_P(LogisticTest, Logistic) {
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
     Value c = logistic(&ctx, a);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, reveal(&ctx, c));
     // low precision
     EXPECT_TRUE(xt::allclose(1.0 / (1.0 + xt::exp(-x)), y, 0.1, 0.5))
         << 1.0 / (1.0 + xt::exp(-x)) << std::endl
diff --git a/libspu/kernel/hal/type_cast_test.cc b/libspu/kernel/hal/type_cast_test.cc
index 247434e8..cd9b4ecc 100644
--- a/libspu/kernel/hal/type_cast_test.cc
+++ b/libspu/kernel/hal/type_cast_test.cc
@@ -32,8 +32,8 @@ TEST(TypeCastTest, int2fxp) {
 
   {
     Value a = constant(&ctx, x, DT_I32);
-    Value c = dtype_cast(&ctx, a, DT_FXP);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    Value c = dtype_cast(&ctx, a, DT_F32);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
     auto y = dump_public_as<float>(&ctx, c);
     EXPECT_TRUE(xt::allclose(expected, y, 0.1, 0.5)) << x << std::endl
@@ -43,10 +43,10 @@ TEST(TypeCastTest, int2fxp) {
 
   {
     Value a = test::makeValue(&ctx, x, VIS_SECRET);
-    Value c = dtype_cast(&ctx, a, DT_FXP);
-    EXPECT_EQ(c.dtype(), DT_FXP);
+    Value c = dtype_cast(&ctx, a, DT_F32);
+    EXPECT_EQ(c.dtype(), DT_F32);
 
-    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).asFxp());
+    auto y = dump_public_as<float>(&ctx, _s2p(&ctx, c).setDtype(DT_F32));
     EXPECT_TRUE(xt::allclose(expected, y, 0.1, 0.5)) << x << std::endl
                                                      << expected << std::endl
                                                      << y;
@@ -62,7 +62,7 @@ TEST(TypeCastTest, fxp2int) {
   SPUContext ctx = test::makeSPUContext();
 
   {
-    Value a = constant(&ctx, x, DT_FXP);
+    Value a = constant(&ctx, x, DT_F32);
     Value c = dtype_cast(&ctx, a, DT_I32);
     EXPECT_EQ(c.dtype(), DT_I32);
 
diff --git a/libspu/kernel/hlo/basic_binary_test.cc b/libspu/kernel/hlo/basic_binary_test.cc
index 6d00f18e..f90c1c01 100644
--- a/libspu/kernel/hlo/basic_binary_test.cc
+++ b/libspu/kernel/hlo/basic_binary_test.cc
@@ -55,7 +55,6 @@ BINARY_EMPTY_TEST(Sub)
 BINARY_EMPTY_TEST(Less)
 BINARY_EMPTY_TEST(Greater)
 BINARY_EMPTY_TEST(Mul)
-BINARY_EMPTY_TEST(Power)
 BINARY_EMPTY_TEST(Max)
 BINARY_EMPTY_TEST(Min)
 BINARY_EMPTY_TEST(And)
@@ -64,6 +63,21 @@ BINARY_EMPTY_TEST(Xor)
 BINARY_EMPTY_TEST(Div)
 BINARY_EMPTY_TEST(Remainder)
 
+TEST_P(BinaryTest, Empty_Power) {
+  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, {0}));
+        auto empty_c1 = Seal(&sctx, Constant(&sctx, 1.0, {0}));
+        auto s_empty = Power(&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);
+      });
+}
+
 TEST_P(BinaryTest, Empty_Dot) {
   FieldType field = std::get<0>(GetParam());
   ProtocolKind prot = std::get<1>(GetParam());
diff --git a/libspu/kernel/hlo/const_test.cc b/libspu/kernel/hlo/const_test.cc
index 17220169..cbcf1181 100644
--- a/libspu/kernel/hlo/const_test.cc
+++ b/libspu/kernel/hlo/const_test.cc
@@ -37,7 +37,7 @@ TEST(ConstTest, Empty) {
 TEST(ConstTest, Epsilon) {
   SPUContext sctx = test::makeSPUContext();
 
-  auto eps = Epsilon(&sctx);
+  auto eps = Epsilon(&sctx, DT_F32);
 
   auto v = hal::dump_public_as<float>(&sctx, eps);
 
diff --git a/libspu/kernel/hlo/shuffle_test.cc b/libspu/kernel/hlo/shuffle_test.cc
index 0550ab58..ce1d494b 100644
--- a/libspu/kernel/hlo/shuffle_test.cc
+++ b/libspu/kernel/hlo/shuffle_test.cc
@@ -18,6 +18,7 @@
 #include "xtensor/xio.hpp"
 #include "xtensor/xsort.hpp"
 
+#include "libspu/kernel/hal/type_cast.h"
 #include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hlo {
@@ -27,12 +28,10 @@ TEST(SortTest, Array) {
   xt::xarray<float> x = xt::random::rand<float>({10});
   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());
+  auto ret1_hat = hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, ret1));
 
   spu::Value ret2 = Shuffle(&ctx, x_v, 0)[0];
-  auto ret2_hat =
-      hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, ret2).asFxp());
+  auto ret2_hat = hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, ret2));
 
   EXPECT_TRUE(xt::allclose(xt::sort(x), xt::sort(ret1_hat), 0.01, 0.001))
       << xt::sort(x) << std::endl
@@ -50,12 +49,10 @@ TEST(SortTest, 2D) {
   std::vector<Value> x_v = {test::makeValue(&ctx, x, VIS_SECRET)};
 
   spu::Value ret1 = Shuffle(&ctx, x_v, 1)[0];
-  auto ret1_hat =
-      hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, ret1).asFxp());
+  auto ret1_hat = hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, ret1));
 
   spu::Value ret2 = Shuffle(&ctx, x_v, 1)[0];
-  auto ret2_hat =
-      hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, ret2).asFxp());
+  auto ret2_hat = hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, ret2));
 
   EXPECT_TRUE(xt::allclose(xt::sort(x, 1), xt::sort(ret1_hat, 1), 0.01, 0.001))
       << xt::sort(x, 1) << std::endl
diff --git a/libspu/kernel/hlo/sort_test.cc b/libspu/kernel/hlo/sort_test.cc
index 874a92db..32fb4737 100644
--- a/libspu/kernel/hlo/sort_test.cc
+++ b/libspu/kernel/hlo/sort_test.cc
@@ -21,6 +21,7 @@
 
 #include "libspu/kernel/hal/constants.h"
 #include "libspu/kernel/hal/polymorphic.h"
+#include "libspu/kernel/hal/type_cast.h"
 #include "libspu/kernel/test_util.h"
 
 namespace spu::kernel::hlo {
@@ -42,7 +43,7 @@ TEST(SortTest, Simple) {
   EXPECT_EQ(rets.size(), 1);
 
   auto sorted_x_hat =
-      hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, rets[0]).asFxp());
+      hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, rets[0]));
 
   EXPECT_TRUE(xt::allclose(sorted_x, sorted_x_hat, 0.01, 0.001))
       << sorted_x << std::endl
@@ -66,7 +67,7 @@ TEST(SortTest, SimpleWithNoPadding) {
   EXPECT_EQ(rets.size(), 1);
 
   auto sorted_x_hat =
-      hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, rets[0]).asFxp());
+      hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, rets[0]));
 
   EXPECT_TRUE(xt::allclose(sorted_x, sorted_x_hat, 0.01, 0.001))
       << sorted_x << std::endl
@@ -93,15 +94,15 @@ TEST(SortTest, MultiInputs) {
   EXPECT_EQ(rets.size(), 2);
 
   auto sorted_x1_hat =
-      hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, rets[0]).asFxp());
+      hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, rets[0]));
 
   EXPECT_TRUE(xt::allclose(sorted_x1, sorted_x1_hat, 0.01, 0.001))
       << sorted_x1 << std::endl
       << sorted_x1_hat << std::endl;
 
   // NOTE: Secret sort is unstable, so rets[1] need to be sort before check.
-  auto sorted_x2_hat = xt::sort(
-      hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, rets[1]).asFxp()));
+  auto sorted_x2_hat =
+      xt::sort(hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, rets[1])));
 
   EXPECT_TRUE(xt::allclose(sorted_x2, sorted_x2_hat, 0.01, 0.001))
       << sorted_x2 << std::endl
@@ -133,9 +134,9 @@ TEST(SortTest, MultiOperands) {
   EXPECT_EQ(rets.size(), 2);
 
   auto sorted_k1_hat =
-      hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, rets[0]).asFxp());
+      hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, rets[0]));
   auto sorted_k2_hat =
-      hal::dump_public_as<float>(&ctx, hal::_s2p(&ctx, rets[1]).asFxp());
+      hal::dump_public_as<float>(&ctx, hal::reveal(&ctx, rets[1]));
 
   EXPECT_TRUE(xt::allclose(sorted_k1, sorted_k1_hat, 0.01, 0.001))
       << sorted_k1 << std::endl
diff --git a/libspu/mpc/cheetah/arith/cheetah_dot_test.cc b/libspu/mpc/cheetah/arith/cheetah_dot_test.cc
index 0a8f2829..ef5fe120 100644
--- a/libspu/mpc/cheetah/arith/cheetah_dot_test.cc
+++ b/libspu/mpc/cheetah/arith/cheetah_dot_test.cc
@@ -76,52 +76,4 @@ TEST_P(CheetahDotTest, Basic) {
   }
 }
 
-TEST_P(CheetahDotTest, Fork) {
-  size_t kWorldSize = 2;
-  auto field = std::get<0>(GetParam());
-  auto dim3 = std::get<1>(GetParam());
-
-  std::vector<ArrayRef> mat(kWorldSize);
-  mat[0] = ring_rand(field, dim3[0] * dim3[1]);
-  mat[1] = ring_rand(field, dim3[1] * dim3[2]);
-
-  std::vector<ArrayRef> result0(kWorldSize);
-  std::vector<ArrayRef> result1(kWorldSize);
-  std::vector<ArrayRef> result2(kWorldSize);
-
-  utils::simulate(kWorldSize, [&](std::shared_ptr<yacl::link::Context> lctx) {
-    int rank = lctx->Rank();
-    bool lhs = rank == 0;
-    auto dot = std::make_shared<CheetahDot>(lctx);
-    auto fork0 = dot->Fork();
-    result0[rank] = dot->DotOLE(mat[rank], dim3, lhs);
-    result1[rank] = fork0->DotOLE(mat[rank], dim3, lhs);
-
-    auto fork1 = dot->Fork();
-    result2[rank] = fork1->DotOLE(mat[rank], dim3, lhs);
-  });
-
-  auto expected = ring_mmul(mat[0], mat[1], dim3[0], dim3[2], dim3[1]);
-  auto computed0 = ring_add(result0[0], result0[1]);
-  auto computed1 = ring_add(result1[0], result1[1]);
-  auto computed2 = ring_add(result2[0], result2[1]);
-  EXPECT_EQ(expected.numel(), computed0.numel());
-  EXPECT_EQ(expected.numel(), computed1.numel());
-  EXPECT_EQ(expected.numel(), computed2.numel());
-
-  const int64_t kMaxDiff = 1;
-  DISPATCH_ALL_FIELDS(field, "_", [&]() {
-    auto e = ArrayView<ring2k_t>(expected);
-    auto c0 = ArrayView<ring2k_t>(computed0);
-    auto c1 = ArrayView<ring2k_t>(computed1);
-    auto c2 = ArrayView<ring2k_t>(computed2);
-
-    for (auto idx = 0; idx < expected.numel(); idx++) {
-      EXPECT_NEAR(e[idx], c0[idx], kMaxDiff);
-      EXPECT_NEAR(e[idx], c1[idx], kMaxDiff);
-      EXPECT_NEAR(e[idx], c2[idx], kMaxDiff);
-    }
-  });
-}
-
 }  // namespace spu::mpc::cheetah
diff --git a/libspu/mpc/cheetah/arith/cheetah_mul_test.cc b/libspu/mpc/cheetah/arith/cheetah_mul_test.cc
index 5bc09e93..39c0232a 100644
--- a/libspu/mpc/cheetah/arith/cheetah_mul_test.cc
+++ b/libspu/mpc/cheetah/arith/cheetah_mul_test.cc
@@ -134,4 +134,99 @@ TEST_P(CheetahMulTest, BasicBinary) {
   });
 }
 
+TEST_P(CheetahMulTest, MixedRingSizeMul) {
+  size_t kWorldSize = 2;
+  auto field = std::get<0>(GetParam());
+  auto n = std::get<1>(GetParam());
+  // Compute Mul on field then on field2
+  FieldType field2;
+  if (field == FM32) {
+    field2 = FM64;
+  } else if (field == FM64) {
+    field2 = FM128;
+  } else {
+    field2 = FM32;
+    std::swap(field, field2);
+  }
+
+  ArrayRef a_bits = ring_rand(field, n);
+  ArrayRef b_bits = ring_rand(field, n);
+  ArrayRef c_bits = ring_rand(field2, n);
+  ArrayRef d_bits = ring_rand(field2, n);
+
+  std::vector<ArrayRef> a_shr(kWorldSize);
+  std::vector<ArrayRef> b_shr(kWorldSize);
+  std::vector<ArrayRef> c_shr(kWorldSize);
+  std::vector<ArrayRef> d_shr(kWorldSize);
+
+  a_shr[0] = ring_rand(field, n);
+  b_shr[0] = ring_rand(field, n);
+  a_shr[1] = ring_sub(a_bits, a_shr[0]);
+  b_shr[1] = ring_sub(b_bits, b_shr[0]);
+
+  c_shr[0] = ring_rand(field2, n);
+  d_shr[0] = ring_rand(field2, n);
+  c_shr[1] = ring_sub(c_bits, c_shr[0]);
+  d_shr[1] = ring_sub(d_bits, d_shr[0]);
+
+  std::vector<ArrayRef> result(kWorldSize);
+  std::vector<ArrayRef> result2(kWorldSize);
+  utils::simulate(kWorldSize, [&](std::shared_ptr<yacl::link::Context> lctx) {
+    int rank = lctx->Rank();
+    // (a0 + a1) * (b0 + b1)
+    // a0*b0 + a0*b1 + a1*b0 + a1*b1
+    auto mul = std::make_shared<CheetahMul>(lctx);
+
+    ArrayRef cross0, cross1;
+    if (rank == 0) {
+      cross0 = mul->MulOLE(a_shr[0], true);
+      cross1 = mul->MulOLE(b_shr[0], true);
+    } else {
+      cross0 = mul->MulOLE(b_shr[1], false);
+      cross1 = mul->MulOLE(a_shr[1], false);
+    }
+
+    result[rank] = ring_mul(a_shr[rank], b_shr[rank]);
+    ring_add_(result[rank], cross0);
+    ring_add_(result[rank], cross1);
+
+    if (rank == 0) {
+      cross0 = mul->MulOLE(c_shr[0], true);
+      cross1 = mul->MulOLE(d_shr[0], true);
+    } else {
+      cross1 = mul->MulOLE(d_shr[1], false);
+      cross0 = mul->MulOLE(c_shr[1], false);
+    }
+
+    result2[rank] = ring_mul(c_shr[rank], d_shr[rank]);
+    ring_add_(result2[rank], cross0);
+    ring_add_(result2[rank], cross1);
+  });
+
+  auto expected = ring_mul(a_bits, b_bits);
+  auto computed = ring_add(result[0], result[1]);
+
+  auto expected2 = ring_mul(c_bits, d_bits);
+  auto computed2 = ring_add(result2[0], result2[1]);
+
+  const int64_t kMaxDiff = 0;
+  DISPATCH_ALL_FIELDS(field, "_", [&]() {
+    auto e = ArrayView<ring2k_t>(expected);
+    auto c = ArrayView<ring2k_t>(computed);
+
+    for (auto idx = 0; idx < expected.numel(); idx++) {
+      ASSERT_NEAR(e[idx], c[idx], kMaxDiff);
+    }
+  });
+
+  DISPATCH_ALL_FIELDS(field2, "_", [&]() {
+    auto e = ArrayView<ring2k_t>(expected2);
+    auto c = ArrayView<ring2k_t>(computed2);
+
+    for (auto idx = 0; idx < expected2.numel(); idx++) {
+      ASSERT_NEAR(e[idx], c[idx], kMaxDiff);
+    }
+  });
+}
+
 }  // namespace spu::mpc::cheetah::test
diff --git a/libspu/mpc/cheetah/ot/basic_ot_prot_test.cc b/libspu/mpc/cheetah/ot/basic_ot_prot_test.cc
index 7ad8d4b3..26c4294f 100644
--- a/libspu/mpc/cheetah/ot/basic_ot_prot_test.cc
+++ b/libspu/mpc/cheetah/ot/basic_ot_prot_test.cc
@@ -316,63 +316,6 @@ TEST_P(BasicOTProtTest, Multiplexer) {
   });
 }
 
-TEST_P(BasicOTProtTest, Fork) {
-  size_t kWorldSize = 2;
-  size_t n = 101;
-  FieldType field = GetParam();
-
-  std::array<ArrayRef, 3> triple[2];
-  std::array<ArrayRef, 3> packed_triple[2];
-
-  utils::simulate(kWorldSize, [&](std::shared_ptr<yacl::link::Context> ctx) {
-    auto conn = std::make_shared<Communicator>(ctx);
-    BasicOTProtocols ot_prot(conn);
-    packed_triple[ctx->Rank()] = ot_prot.AndTriple(field, n, 8 * SizeOf(field));
-    auto fork = ot_prot.Fork();
-
-    triple[ctx->Rank()] = fork->AndTriple(field, n, 1);
-  });
-
-  DISPATCH_ALL_FIELDS(field, "", [&]() {
-    {
-      auto a0 = xt_adapt<ring2k_t>(triple[0][0]);
-      auto b0 = xt_adapt<ring2k_t>(triple[0][1]);
-      auto c0 = xt_adapt<ring2k_t>(triple[0][2]);
-
-      auto a1 = xt_adapt<ring2k_t>(triple[1][0]);
-      auto b1 = xt_adapt<ring2k_t>(triple[1][1]);
-      auto c1 = xt_adapt<ring2k_t>(triple[1][2]);
-
-      for (size_t i = 0; i < n; ++i) {
-        EXPECT_TRUE(a0[i] < 2 && a1[i] < 2);
-        EXPECT_TRUE(b0[i] < 2 && b1[i] < 2);
-        EXPECT_TRUE(c0[i] < 2 && c1[i] < 2);
-
-        ring2k_t e = (a0[i] ^ a1[i]) & (b0[i] ^ b1[i]);
-        ring2k_t c = (c0[i] ^ c1[i]);
-        EXPECT_EQ(e, c);
-      }
-    }
-    {
-      auto a0 = xt_adapt<ring2k_t>(packed_triple[0][0]);
-      auto b0 = xt_adapt<ring2k_t>(packed_triple[0][1]);
-      auto c0 = xt_adapt<ring2k_t>(packed_triple[0][2]);
-      auto a1 = xt_adapt<ring2k_t>(packed_triple[1][0]);
-      auto b1 = xt_adapt<ring2k_t>(packed_triple[1][1]);
-      auto c1 = xt_adapt<ring2k_t>(packed_triple[1][2]);
-
-      size_t nn = a0.size();
-      EXPECT_TRUE(nn * 8 * SizeOf(field) >= n);
-
-      for (size_t i = 0; i < nn; ++i) {
-        ring2k_t e = (a0[i] ^ a1[i]) & (b0[i] ^ b1[i]);
-        ring2k_t c = (c0[i] ^ c1[i]);
-        EXPECT_EQ(e, c);
-      }
-    }
-  });
-}
-
 TEST_P(BasicOTProtTest, CorrelatedAndTriple) {
   size_t kWorldSize = 2;
   size_t n = 10;
diff --git a/spu/tests/BUILD.bazel b/spu/tests/BUILD.bazel
index 16f4b145..fc21239a 100644
--- a/spu/tests/BUILD.bazel
+++ b/spu/tests/BUILD.bazel
@@ -43,6 +43,19 @@ py_test(
     ],
 )
 
+py_test(
+    name = "jnp_aby3_r128_test_x64",
+    timeout = "long",
+    srcs = ["jnp_aby3_r128_test.py"],
+    env = {
+        "ENABLE_X64_TEST": "1",
+    },
+    main = "jnp_aby3_r128_test.py",
+    deps = [
+        ":jnp_testbase",
+    ],
+)
+
 py_test(
     name = "jnp_aby3_r64_test",
     timeout = "long",
@@ -52,6 +65,19 @@ py_test(
     ],
 )
 
+py_test(
+    name = "jnp_aby3_r64_test_x64",
+    timeout = "long",
+    srcs = ["jnp_aby3_r64_test.py"],
+    env = {
+        "ENABLE_X64_TEST": "1",
+    },
+    main = "jnp_aby3_r64_test.py",
+    deps = [
+        ":jnp_testbase",
+    ],
+)
+
 py_test(
     name = "jnp_cheetah_r64_test",
     timeout = "long",
@@ -61,6 +87,19 @@ py_test(
     ],
 )
 
+py_test(
+    name = "jnp_cheetah_r64_test_x64",
+    timeout = "long",
+    srcs = ["jnp_cheetah_r64_test.py"],
+    env = {
+        "ENABLE_X64_TEST": "1",
+    },
+    main = "jnp_cheetah_r64_test.py",
+    deps = [
+        ":jnp_testbase",
+    ],
+)
+
 py_test(
     name = "jnp_semi2k_r128_test",
     timeout = "long",
@@ -70,6 +109,19 @@ py_test(
     ],
 )
 
+py_test(
+    name = "jnp_semi2k_r128_test_x64",
+    timeout = "long",
+    srcs = ["jnp_semi2k_r128_test.py"],
+    env = {
+        "ENABLE_X64_TEST": "1",
+    },
+    main = "jnp_semi2k_r128_test.py",
+    deps = [
+        ":jnp_testbase",
+    ],
+)
+
 py_test(
     name = "jnp_semi2k_r64_test",
     timeout = "long",
@@ -79,6 +131,19 @@ py_test(
     ],
 )
 
+py_test(
+    name = "jnp_semi2k_r64_test_x64",
+    timeout = "long",
+    srcs = ["jnp_semi2k_r64_test.py"],
+    env = {
+        "ENABLE_X64_TEST": "1",
+    },
+    main = "jnp_semi2k_r64_test.py",
+    deps = [
+        ":jnp_testbase",
+    ],
+)
+
 py_test(
     name = "jnp_ref2k_r64_test",
     timeout = "long",
@@ -88,6 +153,19 @@ py_test(
     ],
 )
 
+py_test(
+    name = "jnp_ref2k_r64_test_x64",
+    timeout = "long",
+    srcs = ["jnp_ref2k_r64_test.py"],
+    env = {
+        "ENABLE_X64_TEST": "1",
+    },
+    main = "jnp_ref2k_r64_test.py",
+    deps = [
+        ":jnp_testbase",
+    ],
+)
+
 py_binary(
     name = "jnp_debug",
     srcs = ["jnp_debug.py"],
diff --git a/spu/tests/jnp_testbase.py b/spu/tests/jnp_testbase.py
index 7621db38..27705389 100644
--- a/spu/tests/jnp_testbase.py
+++ b/spu/tests/jnp_testbase.py
@@ -15,6 +15,7 @@
 
 import collections
 import itertools
+from os import getenv
 from enum import Enum
 from functools import partial
 
@@ -26,10 +27,17 @@
 
 from spu.utils.simulation import sim_jax
 
+if getenv("ENABLE_X64_TEST"):
+    from jax import config
+
+    config.update("jax_enable_x64", True)
+    float_dtypes = [np.float64]
+else:
+    float_dtypes = [np.float32]
+
 all_shapes = [(4,), (2, 3, 4)]
 extra_large_shapes = [(100000000,)]
 
-float_dtypes = [np.float32]
 int32_dtypes = [np.int16, np.uint16, np.int32, np.uint32]
 # Note: less by subtract+msb does not work for int64 when Ring=64
 int64_dtypes = [np.int64, np.uint64]