release v1.4.4

webgpu release support
webgpu support
2023-08-16 23:33:22 +08:00 · 2023-08-16 23:31:04 +08:00 · 2023-08-16 23:07:58 +08:00 · 2023-08-16 21:11:58 +08:00 · 2023-08-16 12:50:22 +08:00 · 2023-08-14 23:37:55 +08:00
41 changed files with 69924 additions and 166 deletions
--- a/.gitattributes
+++ b/.gitattributes
@@ -2,6 +2,7 @@ backend-python/rwkv_pip/** linguist-vendored
 backend-python/wkv_cuda_utils/** linguist-vendored
 backend-python/get-pip.py linguist-vendored
 backend-python/convert_model.py linguist-vendored
+backend-python/convert_safetensors.py linguist-vendored
 backend-python/utils/midi.py linguist-vendored
 build/** linguist-vendored
 finetune/lora/** linguist-vendored
--- a/.github/workflows/release.yml
+++ b/.github/workflows/release.yml
@@ -48,6 +48,11 @@ jobs:
        id: cp310
        with:
          python-version: '3.10'
+      - uses: actions-rs/toolchain@v1
+        with:
+          toolchain: stable
+          override: true
+          target: wasm32-unknown-unknown
      - uses: crazy-max/ghaction-chocolatey@v2
        with:
          args: install upx
@@ -60,6 +65,11 @@ jobs:
          Copy-Item -Path "${{ steps.cp310.outputs.python-path }}/../include" -Destination "py310/include" -Recurse
          Copy-Item -Path "${{ steps.cp310.outputs.python-path }}/../libs" -Destination "py310/libs" -Recurse
          ./py310/python -m pip install cyac==1.7
+          git clone https://github.com/josStorer/ai00_rwkv_server --depth=1
+          cd ai00_rwkv_server
+          cargo build --release
+          mv ./target/release/ai00_server.exe ../backend-rust/webgpu_server.exe
+          cd ..
          go install github.com/wailsapp/wails/v2/cmd/wails@latest
          make
          Rename-Item -Path "build/bin/RWKV-Runner.exe" -NewName "RWKV-Runner_windows_x64.exe"
@@ -76,10 +86,23 @@ jobs:
      - uses: actions/setup-go@v4
        with:
          go-version: '1.20.5'
+      - uses: actions-rs/toolchain@v1
+        with:
+          toolchain: stable
+          override: true
+          target: wasm32-unknown-unknown
      - run: |
          sudo apt-get update
          sudo apt-get install upx
          sudo apt-get install build-essential libgtk-3-dev libwebkit2gtk-4.0-dev
+          git clone https://github.com/josStorer/ai00_rwkv_server --depth=1
+          cd ai00_rwkv_server
+          sudo apt-get install libudev-dev
+          sudo apt-get install libasound2-dev
+          rustup target add x86_64-unknown-linux-gnu
+          cargo build --release --target x86_64-unknown-linux-gnu
+          mv ./target/x86_64-unknown-linux-gnu/release/ai00_server ../backend-rust/webgpu_server
+          cd ..
          go install github.com/wailsapp/wails/v2/cmd/wails@latest
          rm -rf ./backend-python/wkv_cuda_utils
          rm ./backend-python/get-pip.py
@@ -101,7 +124,17 @@ jobs:
      - uses: actions/setup-go@v4
        with:
          go-version: '1.20.5'
+      - uses: actions-rs/toolchain@v1
+        with:
+          toolchain: stable
+          override: true
+          target: wasm32-unknown-unknown
      - run: |
+          git clone https://github.com/josStorer/ai00_rwkv_server --depth=1
+          cd ai00_rwkv_server
+          cargo build --release
+          mv ./target/release/ai00_server ../backend-rust/webgpu_server
+          cd ..
          go install github.com/wailsapp/wails/v2/cmd/wails@latest
          rm -rf ./backend-python/wkv_cuda_utils
          rm ./backend-python/get-pip.py
--- a/.gitignore
+++ b/.gitignore
@@ -5,6 +5,8 @@ __pycache__
 .idea
 .vs
 *.pth
+*.st
+*.safetensors
 *.bin
 /config.json
 /cache.json
--- a/CURRENT_CHANGE.md
+++ b/CURRENT_CHANGE.md
@@ -1,9 +1,10 @@
 ## Changes

- japanese UI
- global penalty
- allow custom user_name and assistant_name (`/chat/completions` API)
- update defaultConfigs
+- webgpu support (AMD, Intel, Nvidia, Apple)
+- add rwkv-cuda-beta support (faster)
+- add misc API (`/models` and `/dashboard/billing/credit_grants`)
+- allow multiple systems
+- allow completions input to be null

 ## Install

--- a/backend-golang/rwkv.go
+++ b/backend-golang/rwkv.go
@@ -10,7 +10,7 @@ import (
 	"strings"
 )

-func (a *App) StartServer(python string, port int, host string) (string, error) {
+func (a *App) StartServer(python string, port int, host string, rwkvBeta bool) (string, error) {
 	var err error
 	if python == "" {
 		python, err = GetPython()
@@ -18,7 +18,19 @@ func (a *App) StartServer(python string, port int, host string) (string, error)
 	if err != nil {
 		return "", err
 	}
-	return Cmd(python, "./backend-python/main.py", strconv.Itoa(port), host)
+	args := []string{python, "./backend-python/main.py"}
+	if rwkvBeta {
+		args = append(args, "--rwkv-beta")
+	}
+	args = append(args, "--port", strconv.Itoa(port), "--host", host)
+	return Cmd(args...)
+}
+
+func (a *App) StartWebGPUServer(port int, host string) (string, error) {
+	args := []string{"./backend-rust/webgpu_server"}
+	args = append(args, "-a", "0", "-t", "backend-rust/assets/rwkv_vocab_v20230424.json",
+		"--port", strconv.Itoa(port), "--ip", host)
+	return Cmd(args...)
 }

 func (a *App) ConvertModel(python string, modelPath string, strategy string, outPath string) (string, error) {
@@ -32,6 +44,17 @@ func (a *App) ConvertModel(python string, modelPath string, strategy string, out
 	return Cmd(python, "./backend-python/convert_model.py", "--in", modelPath, "--out", outPath, "--strategy", strategy)
 }

+func (a *App) ConvertSafetensors(python string, modelPath string, outPath string) (string, error) {
+	var err error
+	if python == "" {
+		python, err = GetPython()
+	}
+	if err != nil {
+		return "", err
+	}
+	return Cmd(python, "./backend-python/convert_safetensors.py", "--input", modelPath, "--output", outPath)
+}
+
 func (a *App) ConvertData(python string, input string, outputPrefix string, vocab string) (string, error) {
 	var err error
 	if python == "" {
--- a/backend-python/convert_safetensors.py
+++ b/backend-python/convert_safetensors.py
@@ -0,0 +1,53 @@
+import json
+import os
+import sys
+import copy
+import torch
+from safetensors.torch import load_file, save_file
+
+import argparse
+
+parser = argparse.ArgumentParser()
+parser.add_argument("--input", type=str, help="Path to input pth model")
+parser.add_argument(
+    "--output",
+    type=str,
+    default="./converted.st",
+    help="Path to output safetensors model",
+)
+args = parser.parse_args()
+
+
+def convert_file(
+    pt_filename: str,
+    sf_filename: str,
+):
+    loaded = torch.load(pt_filename, map_location="cpu")
+    if "state_dict" in loaded:
+        loaded = loaded["state_dict"]
+
+    loaded = {k: v.clone().half() for k, v in loaded.items()}
+    for k, v in loaded.items():
+        print(f"{k}\t{v.shape}\t{v.dtype}")
+
+    # For tensors to be contiguous
+    loaded = {k: v.contiguous() for k, v in loaded.items()}
+
+    dirname = os.path.dirname(sf_filename)
+    os.makedirs(dirname, exist_ok=True)
+    save_file(loaded, sf_filename, metadata={"format": "pt"})
+    reloaded = load_file(sf_filename)
+    for k in loaded:
+        pt_tensor = loaded[k]
+        sf_tensor = reloaded[k]
+        if not torch.equal(pt_tensor, sf_tensor):
+            raise RuntimeError(f"The output tensors do not match for key {k}")
+
+
+if __name__ == "__main__":
+    try:
+        convert_file(args.input, args.output)
+        print(f"Saved to {args.output}")
+    except Exception as e:
+        with open("error.txt", "w") as f:
+            f.write(str(e))
--- a/backend-python/dep_check.py
+++ b/backend-python/dep_check.py
@@ -1,3 +1,4 @@
+import safetensors
 import midi2audio
 import mido
 import lm_dataformat
--- a/backend-python/global_var.py
+++ b/backend-python/global_var.py
@@ -1,5 +1,6 @@
 from enum import Enum, auto

+Args = "args"
 Model = "model"
 Model_Status = "model_status"
 Model_Config = "model_config"
--- a/backend-python/main.py
+++ b/backend-python/main.py
@@ -1,5 +1,11 @@
+import time
+
+start_time = time.time()
+
 import os
 import sys
+import argparse
+from typing import Sequence

 sys.path.append(os.path.dirname(os.path.realpath(__file__)))

@@ -12,7 +18,7 @@ from utils.rwkv import *
 from utils.torch import *
 from utils.ngrok import *
 from utils.log import log_middleware
-from routes import completion, config, state_cache, midi
+from routes import completion, config, state_cache, midi, misc
 import global_var

 app = FastAPI(dependencies=[Depends(log_middleware)])
@@ -28,12 +34,18 @@ app.add_middleware(
 app.include_router(completion.router)
 app.include_router(config.router)
 app.include_router(midi.router)
+app.include_router(misc.router)
 app.include_router(state_cache.router)


@app.on_event("startup")
 def init():
    global_var.init()
+    cmd_params = os.environ["RWKV_RUNNER_PARAMS"]
+    global_var.set(
+        global_var.Args, get_args(cmd_params.split(" ") if cmd_params else None)
+    )
+
    state_cache.init()

    set_torch()
@@ -56,9 +68,34 @@ def exit():
    parent.kill()


-if __name__ == "__main__":
-    uvicorn.run(
-        "main:app",
-        port=8000 if len(sys.argv) < 2 else int(sys.argv[1]),
-        host="127.0.0.1" if len(sys.argv) < 3 else sys.argv[2],
+def get_args(args: Union[Sequence[str], None] = None):
+    parser = argparse.ArgumentParser()
+    group = parser.add_argument_group(title="server arguments")
+    group.add_argument(
+        "--port",
+        type=int,
+        default=8000,
+        help="port to run the server on (default: 8000)",
    )
+    group.add_argument(
+        "--host",
+        type=str,
+        default="127.0.0.1",
+        help="host to run the server on (default: 127.0.0.1)",
+    )
+    group = parser.add_argument_group(title="mode arguments")
+    group.add_argument(
+        "--rwkv-beta",
+        action="store_true",
+        help="whether to use rwkv-beta (default: False)",
+    )
+    args = parser.parse_args(args)
+
+    return args
+
+
+if __name__ == "__main__":
+    args = get_args()
+    os.environ["RWKV_RUNNER_PARAMS"] = " ".join(sys.argv[1:])
+    print("--- %s seconds ---" % (time.time() - start_time))
+    uvicorn.run("main:app", port=args.port, host=args.host, workers=1)
--- a/backend-python/requirements.txt
+++ b/backend-python/requirements.txt
--- a/backend-python/requirements_versions.txt
+++ b/backend-python/requirements_versions.txt
--- a/backend-python/requirements_without_cyac.txt
+++ b/backend-python/requirements_without_cyac.txt
--- a/backend-python/routes/completion.py
+++ b/backend-python/routes/completion.py
@@ -2,11 +2,12 @@ import asyncio
 import json
 from threading import Lock
 from typing import List, Union
+from enum import Enum
 import base64

 from fastapi import APIRouter, Request, status, HTTPException
 from sse_starlette.sse import EventSourceResponse
-from pydantic import BaseModel
+from pydantic import BaseModel, Field
 import numpy as np
 import tiktoken
 from utils.rwkv import *
@@ -16,29 +17,35 @@ import global_var
 router = APIRouter()


+class Role(Enum):
+    User = "user"
+    Assistant = "assistant"
+    System = "system"
+
+
 class Message(BaseModel):
-    role: str
-    content: str
+    role: Role
+    content: str = Field(min_length=1)


 class ChatCompletionBody(ModelConfigBody):
-    messages: List[Message]
+    messages: Union[List[Message], None]
    model: str = "rwkv"
    stream: bool = False
-    stop: Union[str, List[str]] = [
+    stop: Union[str, List[str], None] = [
        "\n\nUser",
        "\n\nQuestion",
        "\n\nQ",
        "\n\nHuman",
        "\n\nBob",
    ]
-    user_name: str = None
-    assistant_name: str = None
+    user_name: Union[str, None] = None
+    assistant_name: Union[str, None] = None

    class Config:
        schema_extra = {
            "example": {
-                "messages": [{"role": "user", "content": "hello"}],
+                "messages": [{"role": Role.User.value, "content": "hello"}],
                "model": "rwkv",
                "stream": False,
                "stop": None,
@@ -54,10 +61,10 @@ class ChatCompletionBody(ModelConfigBody):


 class CompletionBody(ModelConfigBody):
-    prompt: Union[str, List[str]]
+    prompt: Union[str, List[str], None]
    model: str = "rwkv"
    stream: bool = False
-    stop: Union[str, List[str]] = None
+    stop: Union[str, List[str], None] = None

    class Config:
        schema_extra = {
@@ -87,7 +94,7 @@ async def eval_rwkv(
    body: ModelConfigBody,
    prompt: str,
    stream: bool,
-    stop: Union[str, List[str]],
+    stop: Union[str, List[str], None],
    chat_mode: bool,
 ):
    global requests_num
@@ -200,7 +207,7 @@ async def eval_rwkv(
                    "choices": [
                        {
                            "message": {
-                                "role": "assistant",
+                                "role": Role.Assistant.value,
                                "content": response,
                            },
                            "index": 0,
@@ -223,17 +230,12 @@ async def chat_completions(body: ChatCompletionBody, request: Request):
    if model is None:
        raise HTTPException(status.HTTP_400_BAD_REQUEST, "model not loaded")

-    question = body.messages[-1]
-    if question.role == "user":
-        question = question.content
-    elif question.role == "system":
-        question = body.messages[-2]
-        if question.role == "user":
-            question = question.content
-        else:
-            raise HTTPException(status.HTTP_400_BAD_REQUEST, "no question found")
-    else:
-        raise HTTPException(status.HTTP_400_BAD_REQUEST, "no question found")
+    if body.messages is None or body.messages == []:
+        raise HTTPException(status.HTTP_400_BAD_REQUEST, "messages not found")
+
+    basic_system: str = ""
+    if body.messages[0].role == Role.System:
+        basic_system = body.messages[0].content

    interface = model.interface
    user = model.user if body.user_name is None else body.user_name
@@ -241,60 +243,61 @@ async def chat_completions(body: ChatCompletionBody, request: Request):

    is_raven = model.rwkv_type == RWKVType.Raven

-    completion_text = (
-        f"""
+    completion_text: str = ""
+    if basic_system == "":
+        completion_text = (
+            f"""
 The following is a coherent verbose detailed conversation between a girl named {bot} and her friend {user}. \
 {bot} is very intelligent, creative and friendly. \
 {bot} is unlikely to disagree with {user}, and {bot} doesn't like to ask {user} questions. \
 {bot} likes to tell {user} a lot about herself and her opinions. \
 {bot} usually gives {user} kind, helpful and informative advices.\n
 """
-        if is_raven
-        else f"{user}{interface} hi\n\n{bot}{interface} Hi. "
-        + "I am your assistant and I will provide expert full response in full details. Please feel free to ask any question and I will always answer it.\n\n"
-    )
-    for message in body.messages:
-        if message.role == "system":
-            completion_text = (
+            if is_raven
+            else (
+                f"{user}{interface} hi\n\n{bot}{interface} Hi. "
+                + "I am your assistant and I will provide expert full response in full details. Please feel free to ask any question and I will always answer it.\n\n"
+            )
+        )
+    elif basic_system != "":
+        completion_text = (
+            (
                f"The following is a coherent verbose detailed conversation between a girl named {bot} and her friend {user}. "
                if is_raven
                else f"{user}{interface} hi\n\n{bot}{interface} Hi. "
-                + message.content.replace("\\n", "\n")
-                .replace("\r\n", "\n")
-                .replace("\n\n", "\n")
-                .replace("\n", " ")
-                .strip()
-                .replace("You are", f"{bot} is" if is_raven else "I am")
-                .replace("you are", f"{bot} is" if is_raven else "I am")
-                .replace("You're", f"{bot} is" if is_raven else "I'm")
-                .replace("you're", f"{bot} is" if is_raven else "I'm")
-                .replace("You", f"{bot}" if is_raven else "I")
-                .replace("you", f"{bot}" if is_raven else "I")
-                .replace("Your", f"{bot}'s" if is_raven else "My")
-                .replace("your", f"{bot}'s" if is_raven else "my")
-                .replace("你", f"{bot}" if is_raven else "我")
-                + "\n\n"
-            )
-            break
-    for message in body.messages:
-        if message.role == "user":
-            completion_text += (
-                f"{user}{interface} "
-                + message.content.replace("\\n", "\n")
-                .replace("\r\n", "\n")
-                .replace("\n\n", "\n")
-                .strip()
-                + "\n\n"
-            )
-        elif message.role == "assistant":
-            completion_text += (
-                f"{bot}{interface} "
-                + message.content.replace("\\n", "\n")
-                .replace("\r\n", "\n")
-                .replace("\n\n", "\n")
-                .strip()
-                + "\n\n"
            )
+            + basic_system.replace("\r\n", "\n")
+            .replace("\r", "\n")
+            .replace("\n\n", "\n")
+            .replace("\n", " ")
+            .strip()
+            .replace("You are", f"{bot} is" if is_raven else "I am")
+            .replace("you are", f"{bot} is" if is_raven else "I am")
+            .replace("You're", f"{bot} is" if is_raven else "I'm")
+            .replace("you're", f"{bot} is" if is_raven else "I'm")
+            .replace("You", f"{bot}" if is_raven else "I")
+            .replace("you", f"{bot}" if is_raven else "I")
+            .replace("Your", f"{bot}'s" if is_raven else "My")
+            .replace("your", f"{bot}'s" if is_raven else "my")
+            .replace("你", f"{bot}" if is_raven else "我")
+            + "\n\n"
+        )
+
+    for message in body.messages[(0 if basic_system == "" else 1) :]:
+        append_message: str = ""
+        if message.role == Role.User:
+            append_message = f"{user}{interface} " + message.content
+        elif message.role == Role.Assistant:
+            append_message = f"{bot}{interface} " + message.content
+        elif message.role == Role.System:
+            append_message = message.content
+        completion_text += (
+            append_message.replace("\r\n", "\n")
+            .replace("\r", "\n")
+            .replace("\n\n", "\n")
+            .strip()
+            + "\n\n"
+        )
    completion_text += f"{bot}{interface}"

    if type(body.stop) == str:
@@ -345,7 +348,7 @@ async def completions(body: CompletionBody, request: Request):


 class EmbeddingsBody(BaseModel):
-    input: Union[str, List[str], List[List[int]]]
+    input: Union[str, List[str], List[List[int]], None]
    model: str = "rwkv"
    encoding_format: str = None
    fast_mode: bool = False
--- a/backend-python/routes/misc.py
+++ b/backend-python/routes/misc.py
@@ -0,0 +1,131 @@
+from fastapi import APIRouter, HTTPException, status
+from utils.rwkv import AbstractRWKV
+import global_var
+
+router = APIRouter()
+
+
+@router.get("/dashboard/billing/credit_grants", tags=["MISC"])
+def credit_grants():
+    return {
+        "object": "credit_summary",
+        "total_granted": 10000,
+        "total_used": 0,
+        "total_available": 10000,
+        "grants": {
+            "object": "list",
+            "data": [
+                {
+                    "object": "credit_grant",
+                    "grant_amount": 10000,
+                    "used_amount": 0,
+                    "effective_at": 1672531200,
+                    "expires_at": 33229440000,
+                }
+            ],
+        },
+    }
+
+
+fake_models = [
+    {
+        "id": "gpt-3.5-turbo",
+        "object": "model",
+        "created": 1677610602,
+        "owned_by": "openai",
+        "permission": [
+            {
+                "id": "modelperm-zy5TOjnE2zVaicIcKO9bQDgX",
+                "object": "model_permission",
+                "created": 1690864883,
+                "allow_create_engine": False,
+                "allow_sampling": True,
+                "allow_logprobs": True,
+                "allow_search_indices": False,
+                "allow_view": True,
+                "allow_fine_tuning": False,
+                "organization": "*",
+                "group": None,
+                "is_blocking": False,
+            }
+        ],
+        "root": "gpt-3.5-turbo",
+        "parent": None,
+    },
+    {
+        "id": "text-davinci-003",
+        "object": "model",
+        "created": 1669599635,
+        "owned_by": "openai-internal",
+        "permission": [
+            {
+                "id": "modelperm-a6niqBmW2JaGmo0fDO7FEt1n",
+                "object": "model_permission",
+                "created": 1690930172,
+                "allow_create_engine": False,
+                "allow_sampling": True,
+                "allow_logprobs": True,
+                "allow_search_indices": False,
+                "allow_view": True,
+                "allow_fine_tuning": False,
+                "organization": "*",
+                "group": None,
+                "is_blocking": False,
+            }
+        ],
+        "root": "text-davinci-003",
+        "parent": None,
+    },
+]
+
+
+@router.get("/v1/models", tags=["MISC"])
+@router.get("/models", tags=["MISC"])
+def models():
+    model: AbstractRWKV = global_var.get(global_var.Model)
+    model_name = model.name if model else "rwkv"
+
+    return {
+        "object": "list",
+        "data": [
+            {
+                "id": model_name,
+                "object": "model",
+                "owned_by": "rwkv",
+                "root": model_name,
+                "parent": None,
+            },
+            *fake_models,
+        ],
+    }
+
+
+@router.get("/v1/models/{model_id}", tags=["MISC"])
+@router.get("/models/{model_id}", tags=["MISC"])
+def model(model_id: str):
+    for fake_model in fake_models:
+        if fake_model["id"] == model_id:
+            return fake_model
+
+    if "rwkv" in model_id.lower():
+        model: AbstractRWKV = global_var.get(global_var.Model)
+        model_name = model.name if model else "rwkv"
+        return {
+            "id": model_name,
+            "object": "model",
+            "owned_by": "rwkv",
+            "root": model_name,
+            "parent": None,
+        }
+
+    raise HTTPException(
+        status.HTTP_404_NOT_FOUND,
+        {
+            "error": {
+                "message": f"The model '{model_id}' does not exist",
+                "type": "invalid_request_error",
+                "param": "model",
+                "code": "model_not_found",
+            }
+        },
+    )
--- a/backend-python/routes/state_cache.py
+++ b/backend-python/routes/state_cache.py
@@ -1,4 +1,4 @@
-from typing import Any, Dict, List
+from typing import Any, Dict, List, Union
 from utils.log import quick_log
 from fastapi import APIRouter, HTTPException, Request, Response, status
 from pydantic import BaseModel
@@ -60,7 +60,7 @@ def enable_state_cache():

 class AddStateBody(BaseModel):
    prompt: str
-    tokens: List[str]
+    tokens: List[Union[str, int]]
    state: Any
    logits: Any

--- a/backend-python/rwkv_pip/beta/cuda/att_one.cu
+++ b/backend-python/rwkv_pip/beta/cuda/att_one.cu
@@ -0,0 +1,124 @@
+#include "ATen/ATen.h"
+#include <cuda_fp16.h>
+#include <cuda_runtime.h>
+#include <torch/extension.h>
+
+#include "element_wise.h"
+#include "util.h"
+
+// Equivalent Python code:
+// ww = t_first + k
+// p = torch.maximum(pp, ww)
+// e1 = torch.exp(pp - p)
+// e2 = torch.exp(ww - p)
+// wkv = ((e1 * aa + e2 * v) / (e1 * bb + e2)).to(dtype=x.dtype)
+// ww = t_decay + pp
+// p = torch.maximum(ww, k)
+// e1 = torch.exp(ww - p)
+// e2 = torch.exp(k - p)
+// t1 = e1 * aa + e2 * v
+// t2 = e1 * bb + e2
+// r = r * wkv
+// return t1, t2, p, r
+struct WkvForwardOne {
+  const float *t_first;
+  const float *k;
+  const float *pp;
+  const float *aa;
+  const float *bb;
+  const float *t_decay;
+  const float *v;
+  /* out */ float *t1;
+  /* out */ float *t2;
+  /* out */ float *p;
+  /* in & out */ half *r;
+
+  __device__ void operator()(int i) const {
+    float ww = t_first[i] + k[i];
+    float pp_ = pp[i];
+    float p_ = (pp_ > ww) ? pp_ : ww;
+    float e1 = expf(pp_ - p_);
+    float e2 = expf(ww - p_);
+    float aa_ = aa[i];
+    float bb_ = bb[i];
+    float v_ = v[i];
+    r[i] = __hmul(r[i], __float2half(((e1 * aa_ + e2 * v_) / (e1 * bb_ + e2))));
+    ww = t_decay[i] + pp_;
+    float k_ = k[i];
+    p_ = (ww > k_) ? ww : k_;
+    e1 = expf(ww - p_);
+    e2 = expf(k_ - p_);
+    t1[i] = e1 * aa_ + e2 * v_;
+    t2[i] = e1 * bb_ + e2;
+    p[i] = p_;
+  }
+};
+
+/*
+   Equivalent Python code:
+   kx = xx * k_mix + sx * (1 - k_mix)
+   vx = xx * v_mix + sx * (1 - v_mix)
+   rx = xx * r_mix + sx * (1 - r_mix)
+*/
+
+struct Mix {
+  const half *xx;
+  const half *sx;
+  const half *k_mix;
+  const half *v_mix;
+  const half *r_mix;
+  /* out */ half *kx;
+  /* out */ half *vx;
+  /* out */ half *rx;
+
+  __device__ void operator()(int i) const {
+    half xx_ = xx[i];
+    half sx_ = sx[i];
+    half k_mix_ = k_mix[i];
+    half v_mix_ = v_mix[i];
+    half r_mix_ = r_mix[i];
+    kx[i] = __hadd(__hmul(xx_, k_mix_),
+                   __hmul(sx_, __hsub(__float2half(1), k_mix_)));
+    vx[i] = __hadd(__hmul(xx_, v_mix_),
+                   __hmul(sx_, __hsub(__float2half(1), v_mix_)));
+    rx[i] = __hadd(__hmul(xx_, r_mix_),
+                   __hmul(sx_, __hsub(__float2half(1), r_mix_)));
+  }
+};
+
+using torch::Tensor;
+
+void gemm_fp16_cublas(Tensor a, Tensor b, Tensor c);
+
+Tensor att_one(Tensor x, Tensor ln_w, Tensor ln_b, Tensor sx, Tensor k_mix,
+             Tensor v_mix, Tensor r_mix, Tensor kw,
+             /* imm */ Tensor kx, Tensor vw, /* imm */ Tensor vx, Tensor rw,
+             /* imm */ Tensor rx, Tensor ow, Tensor t_first,
+             /* imm */ Tensor k, Tensor pp, Tensor ww, Tensor aa, Tensor bb,
+             Tensor t_decay, /* imm */ Tensor v, /* in & out */ Tensor r,
+             /* out */ Tensor x_plus_out, /* out */ Tensor t1,
+             /* out */ Tensor t2, /* out */ Tensor p) {
+  Tensor xx = at::layer_norm(x, {x.size(-1)}, ln_w, ln_b);
+  element_wise(Mix{data_ptr<half>(xx), data_ptr<half>(sx),
+                   data_ptr<half>(k_mix), data_ptr<half>(v_mix),
+                   data_ptr<half>(r_mix), data_ptr<half>(kx),
+                   data_ptr<half>(vx), data_ptr<half>(rx)},
+               x.numel());
+
+  gemm_fp16_cublas(kx, kw, k);
+  gemm_fp16_cublas(vx, vw, v);
+  gemm_fp16_cublas(rx, rw, r);
+  at::sigmoid_(r);
+
+  element_wise(WkvForwardOne{data_ptr<float>(t_first), data_ptr<float>(k),
+                             data_ptr<float>(pp), data_ptr<float>(aa),
+                             data_ptr<float>(bb), data_ptr<float>(t_decay),
+                             data_ptr<float>(v), data_ptr<float>(t1),
+                             data_ptr<float>(t2), data_ptr<float>(p),
+                             data_ptr<half>(r)},
+               x.numel());
+
+  gemm_fp16_cublas(r, ow, x_plus_out);
+  x_plus_out += x;
+  return xx;
+}
--- a/backend-python/rwkv_pip/beta/cuda/att_seq.cu
+++ b/backend-python/rwkv_pip/beta/cuda/att_seq.cu
@@ -0,0 +1,179 @@
+#include "ATen/ATen.h"
+#include <cuda_fp16.h>
+#include <cuda_runtime.h>
+#include <torch/extension.h>
+
+#include "util.h"
+#include "element_wise.h"
+
+using torch::Tensor;
+
+void gemm_fp16_cublas(Tensor a, Tensor b, Tensor c);
+void gemm_fp16_cublas(const void *a, const void *b, void *c, int m,
+                      int n, int k, bool output_fp32);
+
+// based on `kernel_wkv_forward`, fusing more operations
+__global__ void kernel_wkv_forward_new(
+    const int B, const int T, const int C, const float *__restrict__ const _w,
+    const float *__restrict__ const _u, const float *__restrict__ const _k,
+    const float *__restrict__ const _v, const half *__restrict__ const r,
+    half *__restrict__ const _y, float *__restrict__ const _aa,
+    float *__restrict__ const _bb, float *__restrict__ const _pp) {
+  const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  const int _b = idx / C;
+  const int _c = idx % C;
+  const int _offset = _b * T * C + _c;
+  const int _state_offset = _b * C + _c;
+
+  float u = _u[_c];
+  float w = _w[_c];
+  const float *__restrict__ const k = _k + _offset;
+  const float *__restrict__ const v = _v + _offset;
+  half *__restrict__ const y = _y + _offset;
+
+  float aa = _aa[_state_offset];
+  float bb = _bb[_state_offset];
+  float pp = _pp[_state_offset];
+  for (int i = 0; i < T; i++) {
+    const int ii = i * C;
+    const float kk = k[ii];
+    const float vv = v[ii];
+    float ww = u + kk;
+    float p = max(pp, ww);
+    float e1 = exp(pp - p);
+    float e2 = exp(ww - p);
+    y[ii] = __float2half((e1 * aa + e2 * vv) / (e1 * bb + e2));
+    ww = w + pp;
+    p = max(ww, kk);
+    e1 = exp(ww - p);
+    e2 = exp(kk - p);
+    aa = e1 * aa + e2 * vv;
+    bb = e1 * bb + e2;
+    pp = p;
+  }
+  _aa[_state_offset] = aa;
+  _bb[_state_offset] = bb;
+  _pp[_state_offset] = pp;
+}
+
+void cuda_wkv_forward_new(int B, int T, int C, float *w, float *u, float *k,
+                          float *v, half *r, half *y, float *aa, float *bb,
+                          float *pp) {
+  dim3 threadsPerBlock(min(C, 32));
+  assert(B * C % threadsPerBlock.x == 0);
+  dim3 numBlocks(B * C / threadsPerBlock.x);
+  kernel_wkv_forward_new<<<numBlocks, threadsPerBlock>>>(B, T, C, w, u, k, v, r,
+                                                         y, aa, bb, pp);
+}
+
+__global__ void _att_mix(const half *xx, const half *sx, const half *k_mix,
+                        const half *v_mix, const half *r_mix,
+                        const int outer_size, const int inner_size, half *kx,
+                        half *vx, half *rx) {
+  for (int idx2 = blockIdx.x * blockDim.x + threadIdx.x; idx2 < inner_size;
+       idx2 += blockDim.x * gridDim.x) {
+    half k_mix_ = k_mix[idx2];
+    half v_mix_ = v_mix[idx2];
+    half r_mix_ = r_mix[idx2];
+    for (int row = 0; row < outer_size; ++row) {
+      int idx1 = row * inner_size + idx2;
+      half xx_ = xx[idx1];
+      half sx_ = sx[idx1];
+      kx[idx1] = __hadd(__hmul(xx_, k_mix_),
+                        __hmul(sx_, __hsub(__float2half(1), k_mix_)));
+      vx[idx1] = __hadd(__hmul(xx_, v_mix_),
+                        __hmul(sx_, __hsub(__float2half(1), v_mix_)));
+      rx[idx1] = __hadd(__hmul(xx_, r_mix_),
+                        __hmul(sx_, __hsub(__float2half(1), r_mix_)));
+    }
+  }
+}
+
+void att_mix(const half *xx, const half *sx, const half *k_mix,
+            const half *v_mix, const half *r_mix, const int outer_size,
+            const int inner_size, half *kx, half *vx, half *rx) {
+  // 256 is good enough on most GPUs
+  const int32_t BLOCK_SIZE = 256;
+  assert(inner_size % BLOCK_SIZE == 0);
+  _att_mix<<<inner_size / BLOCK_SIZE, BLOCK_SIZE>>>(
+      xx, sx, k_mix, v_mix, r_mix, outer_size, inner_size, kx, vx, rx);
+}
+
+struct InplaceSigmoid {
+  __device__ __forceinline__ half operator()(int i) const {
+    ptr[i] = __float2half(1.0 / (1.0 + exp(-__half2float(ptr[i]))));
+  }
+  half *ptr;
+};
+
+struct InplaceMul {
+  __device__ __forceinline__ half operator()(int i) const {
+    y[i] = __hmul(x[i], y[i]);
+  }
+  half *y;
+  half *x;
+};
+
+/*
+   Equivalent Python code:
+
+   xx = F.layer_norm(x, (x.shape[-1],), weight=ln_w, bias=ln_b)
+   sx = torch.cat((sx.unsqueeze(0), xx[:-1,:]))
+   kx = xx * k_mix + sx * (1 - k_mix)
+   vx = xx * v_mix + sx * (1 - v_mix)
+   rx = xx * r_mix + sx * (1 - r_mix)
+
+   r = torch.sigmoid(gemm(rx, rw))
+   k = gemm(kx, kw, output_dtype=torch.float32)
+   v = gemm(vx, vw, output_dtype=torch.float32)
+
+   T = x.shape[0]
+   for t in range(T):
+       kk = k[t]
+       vv = v[t]
+       ww = t_first + kk
+       p = torch.maximum(pp, ww)
+       e1 = torch.exp(pp - p)
+       e2 = torch.exp(ww - p)
+       sx[t] = ((e1 * aa + e2 * vv) / (e1 * bb + e2)).to(dtype=x.dtype)
+       ww = t_decay + pp
+       p = torch.maximum(ww, kk)
+       e1 = torch.exp(ww - p)
+       e2 = torch.exp(kk - p)
+       aa = e1 * aa + e2 * vv
+       bb = e1 * bb + e2
+       pp = p
+   out = gemm(r * sx, ow)
+   return x + out, xx[-1,:], aa, bb, pp
+*/
+Tensor att_seq(Tensor x, Tensor sx, Tensor ln_w, Tensor ln_b, Tensor k_mix,
+               Tensor v_mix, Tensor r_mix, Tensor kw, Tensor vw, Tensor rw,
+               Tensor ow, Tensor t_first, Tensor pp, Tensor aa, Tensor bb,
+               Tensor t_decay, /* imm */ Tensor buf, /* out */ Tensor x_plus_out) {
+  Tensor xx = at::layer_norm(x, {x.size(-1)}, ln_w, ln_b);
+  sx = at::cat({sx.unsqueeze(0), xx.slice(0, 0, -1)}, 0);
+  char* buf_ptr = (char*)buf.data_ptr();
+  half* kx = (half*)buf_ptr;
+  half* vx = kx + x.numel();
+  half* rx = vx + x.numel();
+  half* wkv_y = rx + x.numel();
+  att_mix(data_ptr<half>(xx), data_ptr<half>(sx), data_ptr<half>(k_mix),
+         data_ptr<half>(v_mix), data_ptr<half>(r_mix), xx.size(0), xx.size(1),
+         kx, vx, rx);
+  float* k = reinterpret_cast<float*>(wkv_y + x.numel());
+  float* v = k + x.size(0) * kw.size(1);
+  half* r = reinterpret_cast<half*>(v + x.size(0) * vw.size(1));
+
+  gemm_fp16_cublas(kx, kw.data_ptr(), k, x.size(0), kw.size(1), kw.size(0), true);
+  gemm_fp16_cublas(vx, vw.data_ptr(), v, x.size(0), vw.size(1), vw.size(0), true);
+  gemm_fp16_cublas(rx, rw.data_ptr(), r, x.size(0), rw.size(1), rw.size(0), false);
+  element_wise(InplaceSigmoid{r}, x.size(0) * rw.size(1));
+  cuda_wkv_forward_new(1, x.size(0), x.size(1), data_ptr<float>(t_decay),
+                       data_ptr<float>(t_first), k, v, r,
+                       wkv_y, data_ptr<float>(aa),
+                       data_ptr<float>(bb), data_ptr<float>(pp));
+  element_wise(InplaceMul{wkv_y, r}, x.numel());
+  gemm_fp16_cublas(wkv_y, ow.data_ptr(), x_plus_out.data_ptr(), x.size(0), ow.size(1), ow.size(0), false);
+  x_plus_out += x;
+  return xx;
+}
--- a/backend-python/rwkv_pip/beta/cuda/element_wise.h
+++ b/backend-python/rwkv_pip/beta/cuda/element_wise.h
@@ -0,0 +1,21 @@
+#include <cassert>
+#include <cstddef>
+#include <cstdint>
+
+template <typename Func> __global__ void _element_wise(Func func, int n) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n;
+       i += blockDim.x * gridDim.x) {
+    func(i);
+  }
+}
+
+// NOTE: packed data type (e.g. float4) is a overkill for current sizes
+// (4096 in 7B model and 768 in 0.1B model),
+// and is not faster than the plain float version.
+template <typename Func>
+void element_wise(Func func, int n) {
+  // 256 is good enough on most GPUs
+  const int32_t BLOCK_SIZE = 256;
+  assert(n % BLOCK_SIZE == 0);
+  _element_wise<<<n / BLOCK_SIZE, BLOCK_SIZE>>>(func, n);
+}
--- a/backend-python/rwkv_pip/beta/cuda/ffn.cu
+++ b/backend-python/rwkv_pip/beta/cuda/ffn.cu
@@ -0,0 +1,165 @@
+#include "ATen/ATen.h"
+#include <cuda_fp16.h>
+#include <cuda_runtime.h>
+#include <torch/extension.h>
+
+#include "element_wise.h"
+#include "util.h"
+
+using torch::Tensor;
+
+void gemm_fp16_cublas(const void *a, const void *b, void *c, int ori_m,
+                      int ori_n, int ori_k, bool output_fp32);
+
+__global__ void _ffn_seq_mix(const half *xx, const half *sx, const half *k_mix,
+                             const half *r_mix, const int outer_size,
+                             const int inner_size, half *kx, half *rx) {
+  for (int idx2 = blockIdx.x * blockDim.x + threadIdx.x; idx2 < inner_size;
+       idx2 += blockDim.x * gridDim.x) {
+    half k_mix_ = k_mix[idx2];
+    half r_mix_ = r_mix[idx2];
+    for (int row = 0; row < outer_size; ++row) {
+      int idx1 = row * inner_size + idx2;
+      half xx_ = xx[idx1];
+      half sx_ = sx[idx1];
+      kx[idx1] = __hadd(__hmul(xx_, k_mix_),
+                        __hmul(sx_, __hsub(__float2half(1), k_mix_)));
+      rx[idx1] = __hadd(__hmul(xx_, r_mix_),
+                        __hmul(sx_, __hsub(__float2half(1), r_mix_)));
+    }
+  }
+}
+
+void ffn_seq_mix(const half *xx, const half *sx, const half *k_mix,
+                 const half *r_mix, const int outer_size, const int inner_size,
+                 half *kx, half *rx) {
+  // 256 is good enough on most GPUs
+  const int32_t BLOCK_SIZE = 256;
+  assert(inner_size % BLOCK_SIZE == 0);
+  _ffn_seq_mix<<<inner_size / BLOCK_SIZE, BLOCK_SIZE>>>(
+      xx, sx, k_mix, r_mix, outer_size, inner_size, kx, rx);
+}
+
+struct InplaceSigmoid {
+  __device__ __forceinline__ void operator()(int i) const {
+    ptr[i] = __float2half(1.0 / (1.0 + exp(-__half2float(ptr[i]))));
+  }
+  half *ptr;
+};
+
+struct InplaceReLUAndSquare {
+  __device__ __forceinline__ void operator()(int i) const {
+    // __hmax is not defined in old cuda
+    if (__hgt(ptr[i], __float2half(0))) {
+      ptr[i] = __hmul(ptr[i], ptr[i]);
+    } else {
+      ptr[i] = __float2half(0);
+    }
+  }
+  half *ptr;
+};
+
+struct InplaceFma {
+  __device__ __forceinline__ void operator()(int i) const {
+    a[i] = __hfma(a[i], b[i], c[i]);
+  }
+  half *a;
+  const half *b;
+  const half *c;
+};
+
+/*
+   Equivalent Python code:
+
+   xx = F.layer_norm(x, (x.shape[-1],), weight=ln_w, bias=ln_b)
+   sx = torch.cat((sx.unsqueeze(0), xx[:-1,:]))
+   kx = xx * k_mix + sx * (1 - k_mix)
+   rx = xx * r_mix + sx * (1 - r_mix)
+
+   r = torch.sigmoid(gemm(rx, rw))
+   vx = torch.square(torch.relu(gemm(kx, kw)))
+   out = r * gemm(vx, vw)
+   return x + out, xx[-1,:]
+*/
+Tensor ffn_seq(Tensor x, Tensor sx, Tensor ln_w, Tensor ln_b, Tensor k_mix,
+               Tensor r_mix, Tensor kw, Tensor vw, Tensor rw,
+               /* imm */ Tensor buf,
+               /* out */ Tensor x_plus_out) {
+  Tensor xx = at::layer_norm(x, {x.size(-1)}, ln_w, ln_b);
+  sx = at::cat({sx.unsqueeze(0), xx.slice(0, 0, -1)}, 0);
+  char *buf_ptr = (char *)buf.data_ptr();
+  half *kx = (half *)buf_ptr;
+  half *rx = kx + x.numel();
+  half *vx = rx + x.numel();
+  half *r = vx + x.size(0) * kw.size(1);
+  ffn_seq_mix(data_ptr<half>(xx), data_ptr<half>(sx), data_ptr<half>(k_mix),
+              data_ptr<half>(r_mix), xx.size(0), xx.size(1), kx, rx);
+
+  gemm_fp16_cublas(rx, rw.data_ptr(), r, x.size(0), rw.size(1), x.size(1),
+                   false);
+  element_wise(InplaceSigmoid{r}, x.size(0) * rw.size(1));
+  gemm_fp16_cublas(kx, kw.data_ptr(), vx, x.size(0), kw.size(1), x.size(1),
+                   false);
+  element_wise(InplaceReLUAndSquare{vx}, x.size(0) * kw.size(1));
+  gemm_fp16_cublas(vx, vw.data_ptr(), x_plus_out.data_ptr(), x.size(0),
+                   vw.size(1), vw.size(0), false);
+  element_wise(InplaceFma{data_ptr<half>(x_plus_out), r, data_ptr<half>(x)},
+               x_plus_out.numel());
+  return xx;
+}
+
+struct FfnOneMix {
+  __device__ __forceinline__ void operator()(int idx) {
+    half k_mix_ = k_mix[idx];
+    half r_mix_ = r_mix[idx];
+    half xx_ = xx[idx];
+    half sx_ = sx[idx];
+    kx[idx] = __hadd(__hmul(xx_, k_mix_),
+                     __hmul(sx_, __hsub(__float2half(1), k_mix_)));
+    rx[idx] = __hadd(__hmul(xx_, r_mix_),
+                     __hmul(sx_, __hsub(__float2half(1), r_mix_)));
+  }
+  half *k_mix;
+  half *r_mix;
+  half *xx;
+  half *sx;
+  half *kx;
+  half *rx;
+};
+
+/*
+  Equivalent Python code:
+
+  xx = F.layer_norm(x, (x.shape[-1],), weight=ln_w, bias=ln_b)
+  kx = xx * k_mix + sx * (1 - k_mix)
+  rx = xx * r_mix + sx * (1 - r_mix)
+
+  r = torch.sigmoid(gemm(rx, rw))
+  vx = torch.square(torch.relu(gemm(kx, kw)))
+  out = r * gemm(vx, vw)
+  return x + out, xx
+*/
+Tensor ffn_one(Tensor x, Tensor sx, Tensor ln_w, Tensor ln_b, Tensor k_mix,
+               Tensor r_mix, Tensor kw, Tensor vw, Tensor rw,
+               /* imm */ Tensor buf,
+               /* out */ Tensor x_plus_out) {
+  Tensor xx = at::layer_norm(x, {x.size(-1)}, ln_w, ln_b);
+  char *buf_ptr = (char *)buf.data_ptr();
+  half *kx = (half *)buf_ptr;
+  half *rx = kx + x.numel();
+  half *vx = rx + x.numel();
+  half *r = vx + x.size(0) * kw.size(1);
+  element_wise(FfnOneMix{data_ptr<half>(k_mix), data_ptr<half>(r_mix),
+                         data_ptr<half>(xx), data_ptr<half>(sx), kx, rx},
+               x.numel());
+  // vector * matrix, so m = 1
+  gemm_fp16_cublas(rx, rw.data_ptr(), r, 1, rw.size(1), rw.size(0), false);
+  element_wise(InplaceSigmoid{r}, rw.size(1));
+  gemm_fp16_cublas(kx, kw.data_ptr(), vx, 1, kw.size(1), kw.size(0), false);
+  element_wise(InplaceReLUAndSquare{vx}, kw.size(1));
+  gemm_fp16_cublas(vx, vw.data_ptr(), x_plus_out.data_ptr(), 1, vw.size(1),
+                   vw.size(0), false);
+  element_wise(InplaceFma{data_ptr<half>(x_plus_out), r, data_ptr<half>(x)},
+               x_plus_out.numel());
+  return xx;
+}
--- a/backend-python/rwkv_pip/beta/cuda/gemm_fp16_cublas.cpp
+++ b/backend-python/rwkv_pip/beta/cuda/gemm_fp16_cublas.cpp
@@ -0,0 +1,80 @@
+#include <cublas_v2.h>
+#include <cuda.h>
+#include <cuda_fp16.h>
+#include <cuda_runtime.h>
+#include <torch/extension.h>
+
+#define CUBLAS_CHECK(condition)                                                \
+  for (cublasStatus_t _cublas_check_status = (condition);                      \
+       _cublas_check_status != CUBLAS_STATUS_SUCCESS;)                         \
+    throw std::runtime_error("cuBLAS error " +                                 \
+                             std::to_string(_cublas_check_status) + " at " +   \
+                             std::to_string(__LINE__));
+
+#define CUDA_CHECK(condition)                                                  \
+  for (cudaError_t _cuda_check_status = (condition);                           \
+       _cuda_check_status != cudaSuccess;)                                     \
+    throw std::runtime_error(                                                  \
+        "CUDA error " + std::string(cudaGetErrorString(_cuda_check_status)) +  \
+        " at " + std::to_string(__LINE__));
+
+cublasHandle_t get_cublas_handle() {
+  static cublasHandle_t cublas_handle = []() {
+    cublasHandle_t handle = nullptr;
+    CUBLAS_CHECK(cublasCreate(&handle));
+#if CUDA_VERSION < 11000
+    CUBLAS_CHECK(cublasSetMathMode(handle, CUBLAS_TENSOR_OP_MATH));
+#else
+    CUBLAS_CHECK(cublasSetMathMode(handle, CUBLAS_DEFAULT_MATH));
+#endif // CUDA_VERSION < 11000
+    return handle;
+  }();
+  return cublas_handle;
+}
+
+/*
+  NOTE: blas gemm is column-major by default, but we need row-major output.
+  The data of row-major, transposed matrix is exactly the same as the
+  column-major, non-transposed matrix, and C = A * B ---> C^T = B^T * A^T
+ */
+void gemm_fp16_cublas(const void *a, const void *b, void *c, int ori_m,
+                      int ori_n, int ori_k, bool output_fp32) {
+  const auto cuda_data_type = CUDA_R_16F;
+  const auto cuda_c_data_type = output_fp32 ? CUDA_R_32F : CUDA_R_16F;
+  const auto compute_type = CUDA_R_32F;
+  const float sp_alpha = 1.f;
+  // use CUBLAS_OP_N. see the notes above
+  const cublasOperation_t cublas_trans_a = CUBLAS_OP_N;
+  const cublasOperation_t cublas_trans_b = CUBLAS_OP_N;
+  // m = (B^T).size(0) = B.size(1) = n;
+  const int cublas_m = ori_n;
+  const int cublas_k = ori_k;
+  // comptiable with rwkv one mode, where 1-D tensor * 2-D tensor
+  // const int n = a.dense_dim() == 1 ? 1 : a.size(0);
+  const int cublas_n = ori_m;
+  const int cublas_lda = cublas_m;
+  const int cublas_ldb = cublas_k;
+  const int cublas_ldc = cublas_m;
+  cublasHandle_t cublas_handle = get_cublas_handle();
+
+#if CUDA_VERSION >= 11000
+  cublasGemmAlgo_t algo = CUBLAS_GEMM_DEFAULT;
+#else
+  cublasGemmAlgo_t algo = CUBLAS_GEMM_DFALT_TENSOR_OP;
+#endif
+  const float sp_beta = 0.f;
+  CUBLAS_CHECK(cublasGemmEx(
+      cublas_handle, cublas_trans_a, cublas_trans_b, cublas_m, cublas_n,
+      cublas_k, &sp_alpha, b, cuda_data_type, cublas_lda,
+      a, cuda_data_type, cublas_ldb, &sp_beta, c,
+      cuda_c_data_type, cublas_ldc, compute_type, algo));
+}
+
+void gemm_fp16_cublas(torch::Tensor a, torch::Tensor b, torch::Tensor c) {
+  // comptiable with rwkv one mode, 1-D tensor * 2-D tensor
+  const int m = a.dense_dim() == 1 ? 1 : a.size(0);
+  const int n = b.size(1);
+  const int k = b.size(0);
+  gemm_fp16_cublas(a.data_ptr(), b.data_ptr(), c.data_ptr(), m, n, k,
+                   c.dtype() == torch::kFloat32);
+}
--- a/backend-python/rwkv_pip/beta/cuda/operators.cu
+++ b/backend-python/rwkv_pip/beta/cuda/operators.cu
@@ -0,0 +1,246 @@
+#include <stdio.h>
+#include <assert.h>
+#include "ATen/ATen.h"
+#include <cuda_fp16.h>
+#define MIN_VALUE (-1e38)
+typedef at::Half fp16;
+__half *cast(fp16 *ptr) {
+    return reinterpret_cast<__half *>(ptr);
+}
+
+template <typename F>
+__global__ void kernel_wkv_forward(const int B, const int T, const int C,
+                               const float *__restrict__ const _w, const float *__restrict__ const _u, const F *__restrict__ const _k, const F *__restrict__ const _v,
+                               F *__restrict__ const _y, float *__restrict__ const _aa, float *__restrict__ const _bb, float *__restrict__ const _pp) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int _b = idx / C;
+    const int _c = idx % C;
+    const int _offset = _b * T * C + _c;
+    const int _state_offset = _b * C + _c;
+
+    float u = _u[_c];
+    float w = _w[_c];
+    const F *__restrict__ const k = _k + _offset;
+    const F *__restrict__ const v = _v + _offset;
+    F *__restrict__ const y = _y + _offset;
+
+    float aa = _aa[_state_offset];
+    float bb = _bb[_state_offset];
+    float pp = _pp[_state_offset];
+    for (int i = 0; i < T; i++) {
+        const int ii = i * C;
+        const float kk = float(k[ii]);
+        const float vv = float(v[ii]);
+        float ww = u + kk;
+        float p = max(pp, ww);
+        float e1 = exp(pp - p);
+        float e2 = exp(ww - p);
+        y[ii] = F((e1 * aa + e2 * vv) / (e1 * bb + e2));
+        ww = w + pp;
+        p = max(ww, kk);
+        e1 = exp(ww - p);
+        e2 = exp(kk - p);
+        aa = e1 * aa + e2 * vv;
+        bb = e1 * bb + e2;
+        pp = p;
+    }
+    _aa[_state_offset] = aa;
+    _bb[_state_offset] = bb;
+    _pp[_state_offset] = pp;
+}
+
+template <typename F>
+void cuda_wkv_forward(int B, int T, int C, float *w, float *u, F *k, F *v, F *y, float *aa, float *bb, float *pp) {
+    dim3 threadsPerBlock( min(C, 32) );
+    assert(B * C % threadsPerBlock.x == 0);
+    dim3 numBlocks(B * C / threadsPerBlock.x);
+    kernel_wkv_forward<<<numBlocks, threadsPerBlock>>>(B, T, C, w, u, k, v, y, aa, bb, pp);
+}
+
+template void cuda_wkv_forward<fp16>(
+    int B, int T, int C,
+    float *w, float *u, fp16 *k, fp16 *v, fp16 *y,
+    float *aa, float *bb, float *pp);
+template void cuda_wkv_forward<float>(
+    int B, int T, int C,
+    float *w, float *u, float *k, float *v, float *y,
+    float *aa, float *bb, float *pp);
+
+__global__ void kernel_mm_seq_fp32i8(
+    const int B, const int N, const int M,
+    const float *__restrict__ const x, const int x_stride,
+    const uint8_t *__restrict__ const w, const int w_stride,
+    const float *__restrict__ const mx,
+    const float *__restrict__ const rx,
+    const float *__restrict__ const my,
+    const float *__restrict__ const ry,
+    float *__restrict__ const y, const int y_stride) {
+
+    const int i = blockIdx.x * blockDim.x + threadIdx.x;
+    const int k = blockIdx.y * blockDim.y + threadIdx.y;
+
+    if (i < B && k < M) {
+        float y_local = 0;
+        for (int j = 0; j < N; ++j) {
+            y_local += x[i * x_stride + j] * (
+                (float(w[j * w_stride + k]) + 0.5f)
+                * rx[k] * ry[j] + mx[k] + my[j]
+            );
+        }
+        y[i * y_stride + k] = y_local;
+    }
+}
+
+template <typename F>
+void cuda_mm8_seq(int B, int N, int M,
+                  F *x, int x_stride,
+                  uint8_t *w, int w_stride,
+                  F *mx, F *rx,
+                  F *my, F *ry,
+                  F *y, int y_stride);
+
+template <>
+void cuda_mm8_seq<float>(int B, int N, int M,
+                         float *x, int x_stride,
+                         uint8_t *w, int w_stride,
+                         float *mx, float *rx,
+                         float *my, float *ry,
+                         float *y, int y_stride) {
+    dim3 blockSize(1, 128);
+    dim3 gridSize((B + blockSize.x - 1) / blockSize.x, (M + blockSize.y - 1) / blockSize.y);
+    kernel_mm_seq_fp32i8<<<gridSize, blockSize>>>(
+        B, N, M, x, x_stride, w, w_stride,
+        mx, rx, my, ry, y, y_stride);
+}
+
+__global__ void kernel_mm_seq_fp16i8(
+    const int B, const int N, const int M,
+    const __half *__restrict__ const x, const int x_stride,
+    const uint8_t *__restrict__ const w, const int w_stride,
+    const __half *__restrict__ const mx,
+    const __half *__restrict__ const rx,
+    const __half *__restrict__ const my,
+    const __half *__restrict__ const ry,
+    __half *__restrict__ const y, const int y_stride) {
+
+    const int i = blockIdx.x * blockDim.x + threadIdx.x;
+    const int k = blockIdx.y * blockDim.y + threadIdx.y;
+
+    if (i < B && k < M) {
+        float y_local = 0;
+        for (int j = 0; j < N; ++j) {
+            y_local += __half2float(x[i * x_stride + j]) * (
+                (float(w[j * w_stride + k]) + 0.5f)
+                * __half2float(rx[k]) * __half2float(ry[j])
+                + __half2float(mx[k]) + __half2float(my[j])
+            );
+        }
+        y[i * y_stride + k] = __float2half(y_local);
+    }
+}
+
+template <>
+void cuda_mm8_seq<fp16>(int B, int N, int M,
+                        fp16 *x, int x_stride,
+                        uint8_t *w, int w_stride,
+                        fp16 *mx, fp16 *rx,
+                        fp16 *my, fp16 *ry,
+                        fp16 *y, int y_stride) {
+    dim3 blockSize(1, 128);
+    dim3 gridSize((B + blockSize.x - 1) / blockSize.x, (M + blockSize.y - 1) / blockSize.y);
+    kernel_mm_seq_fp16i8<<<gridSize, blockSize>>>(
+        B, N, M, cast(x), x_stride, w, w_stride,
+        cast(mx), cast(rx), cast(my), cast(ry), cast(y), y_stride);
+}
+
+#define MM8_ONE_JSPLIT 24
+#define MM8_ONE_TILE 1024
+
+__global__ void kernel_mm_one_fp32i8(
+    const int N, const int M,
+    const float *__restrict__ const x,
+    const uint8_t *__restrict__ const w, const int w_stride,
+    const float *__restrict__ const mx,
+    const float *__restrict__ const rx,
+    const float *__restrict__ const my,
+    const float *__restrict__ const ry,
+    float *__restrict__ const y) {
+
+    const int k = blockIdx.y * blockDim.y + threadIdx.y;
+    const int j0 = min(N, blockIdx.x * ((N + MM8_ONE_JSPLIT - 1) / MM8_ONE_JSPLIT));
+    const int j1 = min(N, (blockIdx.x + 1) * ((N + MM8_ONE_JSPLIT - 1) / MM8_ONE_JSPLIT));
+
+    if (k < M) {
+        float y_local = 0;
+        for (int j = j0; j < j1; ++j) {
+            y_local += x[j] * (
+                (float(w[j * w_stride + k]) + 0.5f)
+                * rx[k] * ry[j] + mx[k] + my[j]
+            );
+        }
+        atomicAdd(&y[k], y_local);
+    }
+}
+
+template <typename F>
+void cuda_mm8_one(int N, int M,
+                  F *x,
+                  uint8_t *w, int w_stride,
+                  F *mx, F *rx,
+                  F *my, F *ry,
+                  float *y);
+
+template <>
+void cuda_mm8_one<float>(int N, int M,
+                        float *x,
+                        uint8_t *w, int w_stride,
+                        float *mx, float *rx,
+                        float *my, float *ry,
+                        float *y) {
+    dim3 blockSize(1, MM8_ONE_TILE);
+    dim3 gridSize(MM8_ONE_JSPLIT, (M + blockSize.y - 1) / blockSize.y);
+    kernel_mm_one_fp32i8<<<gridSize, blockSize>>>(
+        N, M, x, w, w_stride,
+        mx, rx, my, ry, y);
+}
+
+__global__ void kernel_mm_one_fp16i8(
+    const int N, const int M,
+    const __half *__restrict__ const x,
+    const uint8_t *__restrict__ const w, const int w_stride,
+    const __half *__restrict__ const mx,
+    const __half *__restrict__ const rx,
+    const __half *__restrict__ const my,
+    const __half *__restrict__ const ry,
+    float *__restrict__ const y) {
+
+    const int k = blockIdx.y * blockDim.y + threadIdx.y;
+    const int j0 = min(N, blockIdx.x * ((N + MM8_ONE_JSPLIT - 1) / MM8_ONE_JSPLIT));
+    const int j1 = min(N, (blockIdx.x + 1) * ((N + MM8_ONE_JSPLIT - 1) / MM8_ONE_JSPLIT));
+
+    if (k < M) {
+        float y_local = 0;
+        for (int j = j0; j < j1; ++j) {
+            y_local += __half2float(x[j]) * (
+                (float(w[j * w_stride + k]) + 0.5f)
+                * __half2float(rx[k]) * __half2float(ry[j])
+                + __half2float(mx[k]) + __half2float(my[j])
+            );
+        }
+        atomicAdd(&y[k], y_local);
+    }
+}
+
+template <>
+void cuda_mm8_one<fp16>(int N, int M,
+                        fp16 *x,
+                        uint8_t *w, int w_stride,
+                        fp16 *mx, fp16 *rx,
+                        fp16 *my, fp16 *ry,
+                        float *y) {
+    dim3 blockSize(1, MM8_ONE_TILE);
+    dim3 gridSize(MM8_ONE_JSPLIT, (M + blockSize.y - 1) / blockSize.y);
+    kernel_mm_one_fp16i8<<<gridSize, blockSize>>>(
+        N, M, cast(x), w, w_stride,
+        cast(mx), cast(rx), cast(my), cast(ry), y);
+}
--- a/backend-python/rwkv_pip/beta/cuda/util.h
+++ b/backend-python/rwkv_pip/beta/cuda/util.h
@@ -0,0 +1,7 @@
+#include "ATen/ATen.h"
+#include <cuda_fp16.h>
+
+template <typename T> T *data_ptr(torch::Tensor x) { return x.data_ptr<T>(); }
+template <> inline half *data_ptr(torch::Tensor x) {
+  return reinterpret_cast<half *>(x.data_ptr<at::Half>());
+}
--- a/backend-python/rwkv_pip/beta/cuda/wrapper.cpp
+++ b/backend-python/rwkv_pip/beta/cuda/wrapper.cpp
@@ -0,0 +1,167 @@
+#include <torch/extension.h>
+#include "ATen/ATen.h"
+#include <iostream>
+#include <c10/cuda/CUDAGuard.h>
+
+typedef at::Half fp16;
+
+template <typename F>
+void cuda_wkv_forward(int B, int T, int C,
+                      float *w, float *u, F *k, F *v, F *y,
+                      float *aa, float *bb, float *pp);
+template <typename F>
+void cuda_mm8_seq(int B, int N, int M,
+                  F *x, int x_stride,
+                  uint8_t *w, int w_stride,
+                  F *mx, F *rx,
+                  F *my, F *ry,
+                  F *y, int y_stride);
+template <typename F>
+void cuda_mm8_one(int N, int M,
+                  F *x,
+                  uint8_t *w, int w_stride,
+                  F *mx, F *rx,
+                  F *my, F *ry,
+                  float *y);
+
+void wkv_forward(int64_t B, int64_t T, int64_t C,
+                 torch::Tensor &w, torch::Tensor &u,
+                 torch::Tensor &k, torch::Tensor &v, torch::Tensor &y,
+                 torch::Tensor &aa, torch::Tensor &bb, torch::Tensor &pp) {
+    const at::cuda::OptionalCUDAGuard device_guard(device_of(w));
+    switch (k.scalar_type()) {
+    case c10::ScalarType::Half:
+        cuda_wkv_forward(B, T, C,
+                         w.data_ptr<float>(), u.data_ptr<float>(),
+                         k.data_ptr<fp16>(), v.data_ptr<fp16>(), y.data_ptr<fp16>(),
+                         aa.data_ptr<float>(), bb.data_ptr<float>(), pp.data_ptr<float>());
+        break;
+    case c10::ScalarType::Float:
+        cuda_wkv_forward(B, T, C,
+                         w.data_ptr<float>(), u.data_ptr<float>(),
+                         k.data_ptr<float>(), v.data_ptr<float>(), y.data_ptr<float>(),
+                         aa.data_ptr<float>(), bb.data_ptr<float>(), pp.data_ptr<float>());
+        break;
+    default:
+        assert(false && "Only FP16 and FP32 are currently supported");
+    }
+}
+
+void mm8_seq(int64_t B, int64_t N, int64_t M,
+             torch::Tensor &x, torch::Tensor &w,
+             torch::Tensor &mx, torch::Tensor &rx,
+             torch::Tensor &my, torch::Tensor &ry,
+             torch::Tensor &y) {
+    assert(x.stride(1) == 1);
+    assert(w.stride(1) == 1);
+    assert(mx.stride(0) == 1 && rx.stride(0) == 1);
+    assert(my.stride(0) == 1 && ry.stride(0) == 1);
+    assert(y.stride(1) == 1);
+    const at::cuda::OptionalCUDAGuard device_guard(device_of(w));
+    switch (x.scalar_type()) {
+    case c10::ScalarType::Half:
+        cuda_mm8_seq(
+            B, N, M,
+            x.data_ptr<fp16>(), x.stride(0),
+            w.data_ptr<uint8_t>(), w.stride(0),
+            mx.data_ptr<fp16>(), rx.data_ptr<fp16>(),
+            my.data_ptr<fp16>(), ry.data_ptr<fp16>(),
+            y.data_ptr<fp16>(), y.stride(0));
+        break;
+    case c10::ScalarType::Float:
+        cuda_mm8_seq(
+            B, N, M,
+            x.data_ptr<float>(), x.stride(0),
+            w.data_ptr<uint8_t>(), w.stride(0),
+            mx.data_ptr<float>(), rx.data_ptr<float>(),
+            my.data_ptr<float>(), ry.data_ptr<float>(),
+            y.data_ptr<float>(), y.stride(0));
+        break;
+    default:
+        assert(false && "Only FP16 and FP32 are currently supported");
+    }
+}
+void mm8_one(int64_t N, int64_t M,
+             torch::Tensor &x, torch::Tensor &w,
+             torch::Tensor &mx, torch::Tensor &rx,
+             torch::Tensor &my, torch::Tensor &ry,
+             torch::Tensor &y) {
+    assert(x.stride(0) == 1);
+    assert(w.stride(1) == 1);
+    assert(mx.stride(0) == 1 && rx.stride(0) == 1);
+    assert(my.stride(0) == 1 && ry.stride(0) == 1);
+    assert(y.stride(0) == 1);
+    const at::cuda::OptionalCUDAGuard device_guard(device_of(w));
+    switch (x.scalar_type()) {
+    case c10::ScalarType::Half:
+        cuda_mm8_one(
+            N, M,
+            x.data_ptr<fp16>(),
+            w.data_ptr<uint8_t>(), w.stride(0),
+            mx.data_ptr<fp16>(), rx.data_ptr<fp16>(),
+            my.data_ptr<fp16>(), ry.data_ptr<fp16>(),
+            y.data_ptr<float>());
+        break;
+    case c10::ScalarType::Float:
+        cuda_mm8_one(
+            N, M,
+            x.data_ptr<float>(),
+            w.data_ptr<uint8_t>(), w.stride(0),
+            mx.data_ptr<float>(), rx.data_ptr<float>(),
+            my.data_ptr<float>(), ry.data_ptr<float>(),
+            y.data_ptr<float>());
+        break;
+    default:
+        assert(false && "Only FP16 and FP32 are currently supported");
+    }
+}
+
+using torch::Tensor;
+
+void gemm_fp16_cublas(Tensor a, Tensor b, Tensor c);
+
+Tensor att_one(Tensor x, Tensor ln_w, Tensor ln_b, Tensor sx, Tensor k_mix,
+             Tensor v_mix, Tensor r_mix, Tensor kw,
+             /* imm */ Tensor kx, Tensor vw, /* imm */ Tensor vx, Tensor rw,
+             /* imm */ Tensor rx, Tensor ow, Tensor t_first,
+             /* imm */ Tensor k, Tensor pp, Tensor ww, Tensor aa, Tensor bb,
+             Tensor t_decay, /* imm */ Tensor v, /* in & out */ Tensor r,
+             /* out */ Tensor x_plus_out, /* out */ Tensor t1,
+             /* out */ Tensor t2, /* out */ Tensor p);
+
+Tensor att_seq(Tensor x, Tensor sx, Tensor ln_w, Tensor ln_b, Tensor k_mix,
+               Tensor v_mix, Tensor r_mix, Tensor kw, Tensor vw, Tensor rw,
+               Tensor ow, Tensor t_first, Tensor pp, Tensor aa, Tensor bb,
+               Tensor t_decay, /* imm */ Tensor buf, /* out */ Tensor x_plus_out);
+
+Tensor ffn_seq(Tensor x, Tensor sx, Tensor ln_w, Tensor ln_b, Tensor k_mix,
+               Tensor r_mix, Tensor kw, Tensor vw, Tensor rw,
+               /* imm */ Tensor buf,
+               /* out */ Tensor x_plus_out);
+
+Tensor ffn_one(Tensor x, Tensor sx, Tensor ln_w, Tensor ln_b, Tensor k_mix,
+               Tensor r_mix, Tensor kw, Tensor vw, Tensor rw,
+               /* imm */ Tensor buf,
+               /* out */ Tensor x_plus_out);
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+    m.def("wkv_forward", &wkv_forward, "wkv forward");
+    m.def("mm8_seq", &mm8_seq, "mm8 seq");
+    m.def("mm8_one", &mm8_one, "mm8 one");
+    m.def("gemm_fp16_cublas", &gemm_fp16_cublas, "gemv fp16 cublas");
+    m.def("att_one", &att_one, "att one");
+    m.def("att_seq", &att_seq, "att seq");
+    m.def("ffn_seq", &ffn_seq, "ffn seq");
+    m.def("ffn_one", &ffn_one, "ffn one");
+}
+
+TORCH_LIBRARY(rwkv, m) {
+    m.def("wkv_forward", wkv_forward);
+    m.def("mm8_seq", mm8_seq);
+    m.def("mm8_one", mm8_one);
+    m.def("gemm_fp16_cublas", gemm_fp16_cublas);
+    m.def("att_one", att_one);
+    m.def("att_seq", att_seq);
+    m.def("ffn_seq", ffn_seq);
+    m.def("ffn_one", ffn_one);
+}
--- a/backend-python/rwkv_pip/beta/model.py
+++ b/backend-python/rwkv_pip/beta/model.py
--- a/backend-python/utils/log.py
+++ b/backend-python/utils/log.py
@@ -2,6 +2,8 @@ import json
 import logging
 from typing import Any
 from fastapi import Request
+from pydantic import BaseModel
+from enum import Enum


 logger = logging.getLogger()
@@ -14,12 +16,21 @@ fh.setFormatter(formatter)
 logger.addHandler(fh)


+class ClsEncoder(json.JSONEncoder):
+    def default(self, obj):
+        if isinstance(obj, BaseModel):
+            return obj.dict()
+        if isinstance(obj, Enum):
+            return obj.value
+        return super().default(obj)
+
+
 def quick_log(request: Request, body: Any, response: str):
    try:
        logger.info(
            f"Client: {request.client if request else ''}\nUrl: {request.url if request else ''}\n"
            + (
-                f"Body: {json.dumps(body.__dict__, default=vars, ensure_ascii=False)}\n"
+                f"Body: {json.dumps(body.__dict__, ensure_ascii=False, cls=ClsEncoder)}\n"
                if body
                else ""
            )
--- a/backend-python/utils/rwkv.py
+++ b/backend-python/utils/rwkv.py
@@ -10,6 +10,7 @@ from fastapi import HTTPException
 from pydantic import BaseModel, Field
 import numpy as np
 from routes import state_cache
+import global_var


 END_OF_TEXT = 0
@@ -27,7 +28,17 @@ class RWKVType(Enum):

 class AbstractRWKV(ABC):
    def __init__(self, model: str, strategy: str, tokens_path: str):
-        from rwkv.model import RWKV as Model  # dynamic import to make RWKV_CUDA_ON work
+        rwkv_beta = global_var.get(global_var.Args).rwkv_beta
+
+        # dynamic import to make RWKV_CUDA_ON work
+        if rwkv_beta:
+            from rwkv_pip.beta.model import (
+                RWKV as Model,
+            )
+        else:
+            from rwkv.model import (
+                RWKV as Model,
+            )
        from rwkv_pip.utils import PIPELINE

        filename, _ = os.path.splitext(os.path.basename(model))
@@ -221,7 +232,7 @@ class AbstractRWKV(ABC):
                return state[0].tolist(), token_len

    def generate(
-        self, prompt: str, stop: Union[str, List[str]] = None
+        self, prompt: str, stop: Union[str, List[str], None] = None
    ) -> Iterable[Tuple[str, str, int, int]]:
        quick_log(None, None, "Generation Prompt:\n" + prompt)
        cache = None
@@ -438,8 +449,10 @@ The following is a coherent verbose detailed conversation between a girl named {
 {bot} usually gives {user} kind, helpful and informative advices.\n
 """
            if self.rwkv_type == RWKVType.Raven
-            else f"{user}{interface} hi\n\n{bot}{interface} Hi. "
-            + "I am your assistant and I will provide expert full response in full details. Please feel free to ask any question and I will always answer it.\n\n"
+            else (
+                f"{user}{interface} hi\n\n{bot}{interface} Hi. "
+                + "I am your assistant and I will provide expert full response in full details. Please feel free to ask any question and I will always answer it.\n\n"
+            )
        )
        logits, _ = self.run_rnn(self.fix_tokens(self.pipeline.encode(preset_system)))
        try:
--- a/backend-rust/assets/rwkv_vocab_v20230424.json
+++ b/backend-rust/assets/rwkv_vocab_v20230424.json
--- a/frontend/src/_locales/ja/main.json
+++ b/frontend/src/_locales/ja/main.json
@@ -128,6 +128,7 @@
  "Chinese Kongfu": "中国武術",
  "Allow external access to the API (service must be restarted)": "APIへの外部アクセスを許可する (サービスを再起動する必要があります)",
  "Custom": "カスタム",
+  "CUDA (Beta, Faster)": "CUDA (ベータ、高速)",
  "Reset All Configs": "すべての設定をリセット",
  "Cancel": "キャンセル",
  "Confirm": "確認",
@@ -240,5 +241,8 @@
  "Auto Play At The End": "最後に自動再生",
  "No File to save": "保存するファイルがありません",
  "File Saved": "ファイルが保存されました",
-  "Failed to load local sound font, please check if the files exist - assets/sound-font": "ローカルサウンドフォントの読み込みに失敗しました、ファイルが存在するか確認してください - assets/sound-font"
+  "Failed to load local sound font, please check if the files exist - assets/sound-font": "ローカルサウンドフォントの読み込みに失敗しました、ファイルが存在するか確認してください - assets/sound-font",
+  "Please convert model to safe tensors format first": "モデルを安全なテンソル形式に変換してください",
+  "Convert To Safe Tensors Format": "安全なテンソル形式に変換",
+  "Please change Strategy to WebGPU to use safetensors format": "StrategyをWebGPUに変更して、安全なテンソル形式を使用してください"
 }
--- a/frontend/src/_locales/zh-hans/main.json
+++ b/frontend/src/_locales/zh-hans/main.json
@@ -128,6 +128,7 @@
  "Chinese Kongfu": "情境冒险",
  "Allow external access to the API (service must be restarted)": "允许外部访问API (必须重启服务)",
  "Custom": "自定义",
+  "CUDA (Beta, Faster)": "CUDA (Beta, 更快)",
  "Reset All Configs": "重置所有配置",
  "Cancel": "取消",
  "Confirm": "确认",
@@ -240,5 +241,8 @@
  "Auto Play At The End": "结束时自动播放",
  "No File to save": "无文件可保存",
  "File Saved": "文件已保存",
-  "Failed to load local sound font, please check if the files exist - assets/sound-font": "加载本地音色资源失败，请检查文件是否存在 - assets/sound-font"
+  "Failed to load local sound font, please check if the files exist - assets/sound-font": "加载本地音色资源失败，请检查文件是否存在 - assets/sound-font",
+  "Please convert model to safe tensors format first": "请先将模型转换为Safetensors格式",
+  "Convert To Safe Tensors Format": "转换为Safetensors格式",
+  "Please change Strategy to WebGPU to use safetensors format": "请将Strategy改为WebGPU以使用safetensors格式"
 }
--- a/frontend/src/components/ResetConfigsButton.tsx
+++ b/frontend/src/components/ResetConfigsButton.tsx
@@ -11,7 +11,7 @@ export const ResetConfigsButton: FC<{ afterConfirm?: () => void }> = ({ afterCon
  return <DialogButton icon={<ArrowReset20Regular />} tooltip={t('Reset All Configs')} title={t('Reset All Configs')}
    contentText={t('Are you sure you want to reset all configs? This will obtain the latest preset configs, but will override your custom configs and cannot be undone.')}
    onConfirm={() => {
-      commonStore.setModelConfigs(commonStore.platform != 'darwin' ? defaultModelConfigs : defaultModelConfigsMac, false);
+      commonStore.setModelConfigs(commonStore.platform !== 'darwin' ? defaultModelConfigs : defaultModelConfigsMac, false);
      commonStore.setCurrentConfigIndex(0, true);
      afterConfirm?.();
    }} />;
--- a/frontend/src/components/RunButton.tsx
+++ b/frontend/src/components/RunButton.tsx
@@ -1,6 +1,12 @@
 import React, { FC, MouseEventHandler, ReactElement } from 'react';
 import commonStore, { ModelStatus } from '../stores/commonStore';
-import { AddToDownloadList, CopyFile, FileExists, StartServer } from '../../wailsjs/go/backend_golang/App';
+import {
+  AddToDownloadList,
+  CopyFile,
+  FileExists,
+  StartServer,
+  StartWebGPUServer
+} from '../../wailsjs/go/backend_golang/App';
 import { Button } from '@fluentui/react-components';
 import { observer } from 'mobx-react-lite';
 import { exit, getStatus, readRoot, switchModel, updateConfig } from '../apis';
@@ -39,6 +45,7 @@ export const RunButton: FC<{ onClickRun?: MouseEventHandler, iconMode?: boolean
      commonStore.setStatus({ status: ModelStatus.Starting });

      const modelConfig = commonStore.getCurrentModelConfig();
+      const webgpu = modelConfig.modelParameters.device === 'WebGPU';
      let modelName = '';
      let modelPath = '';
      if (modelConfig && modelConfig.modelParameters) {
@@ -50,9 +57,32 @@ export const RunButton: FC<{ onClickRun?: MouseEventHandler, iconMode?: boolean
        return;
      }

-      const ok = await checkDependencies(navigate);
-      if (!ok)
-        return;
+      if (webgpu) {
+        if (!['.st', '.safetensors'].some(ext => modelPath.endsWith(ext))) {
+          const stModelPath = modelPath.replace(/\.pth$/, '.st');
+          if (await FileExists(stModelPath)) {
+            modelPath = stModelPath;
+          } else {
+            toast(t('Please convert model to safe tensors format first'), { type: 'error' });
+            commonStore.setStatus({ status: ModelStatus.Offline });
+            return;
+          }
+        }
+      }
+
+      if (!webgpu) {
+        if (['.st', '.safetensors'].some(ext => modelPath.endsWith(ext))) {
+          toast(t('Please change Strategy to WebGPU to use safetensors format'), { type: 'error' });
+          commonStore.setStatus({ status: ModelStatus.Offline });
+          return;
+        }
+      }
+
+      if (!webgpu) {
+        const ok = await checkDependencies(navigate);
+        if (!ok)
+          return;
+      }

      const currentModelSource = commonStore.modelSourceList.find(item => item.name === modelName);

@@ -85,7 +115,14 @@ export const RunButton: FC<{ onClickRun?: MouseEventHandler, iconMode?: boolean

      await exit(1000).catch(() => {
      });
-      StartServer(commonStore.settings.customPythonPath, port, commonStore.settings.host !== '127.0.0.1' ? '0.0.0.0' : '127.0.0.1').catch((e) => {
+
+      const startServer = webgpu ?
+        (_: string, port: number, host: string) => StartWebGPUServer(port, host)
+        : StartServer;
+
+      startServer(commonStore.settings.customPythonPath, port, commonStore.settings.host !== '127.0.0.1' ? '0.0.0.0' : '127.0.0.1',
+        modelConfig.modelParameters.device === 'CUDA-Beta'
+      ).catch((e) => {
        const errMsg = e.message || e;
        if (errMsg.includes('path contains space'))
          toast(`${t('Error')} - ${t('File Path Cannot Contain Space')}`, { type: 'error' });
@@ -102,23 +139,27 @@ export const RunButton: FC<{ onClickRun?: MouseEventHandler, iconMode?: boolean
          if (r.ok && !loading) {
            loading = true;
            clearInterval(intervalId);
-            await getStatus().then(status => {
-              if (status)
-                commonStore.setStatus(status);
-            });
+            if (!webgpu) {
+              await getStatus().then(status => {
+                if (status)
+                  commonStore.setStatus(status);
+              });
+            }
            commonStore.setStatus({ status: ModelStatus.Loading });
            toast(t('Loading Model'), { type: 'info' });
-            updateConfig({
-              max_tokens: modelConfig.apiParameters.maxResponseToken,
-              temperature: modelConfig.apiParameters.temperature,
-              top_p: modelConfig.apiParameters.topP,
-              presence_penalty: modelConfig.apiParameters.presencePenalty,
-              frequency_penalty: modelConfig.apiParameters.frequencyPenalty
-            });
+            if (!webgpu) {
+              updateConfig({
+                max_tokens: modelConfig.apiParameters.maxResponseToken,
+                temperature: modelConfig.apiParameters.temperature,
+                top_p: modelConfig.apiParameters.topP,
+                presence_penalty: modelConfig.apiParameters.presencePenalty,
+                frequency_penalty: modelConfig.apiParameters.frequencyPenalty
+              });
+            }

            const strategy = getStrategy(modelConfig);
            let customCudaFile = '';
-            if ((modelConfig.modelParameters.device === 'CUDA' || modelConfig.modelParameters.device === 'Custom')
+            if ((modelConfig.modelParameters.device.includes('CUDA') || modelConfig.modelParameters.device === 'Custom')
              && modelConfig.modelParameters.useCustomCuda && !strategy.includes('fp32')) {
              if (commonStore.platform === 'windows') {
                customCudaFile = getSupportedCustomCudaFile();
--- a/frontend/src/pages/Configs.tsx
+++ b/frontend/src/pages/Configs.tsx
@@ -13,13 +13,14 @@ import { Page } from '../components/Page';
 import { useNavigate } from 'react-router';
 import { RunButton } from '../components/RunButton';
 import { updateConfig } from '../apis';
-import { ConvertModel, FileExists, GetPyError } from '../../wailsjs/go/backend_golang/App';
-import { getStrategy } from '../utils';
+import { ConvertModel, ConvertSafetensors, FileExists, GetPyError } from '../../wailsjs/go/backend_golang/App';
+import { checkDependencies, getStrategy } from '../utils';
 import { useTranslation } from 'react-i18next';
 import { WindowShow } from '../../wailsjs/runtime/runtime';
 import strategyImg from '../assets/images/strategy.jpg';
 import strategyZhImg from '../assets/images/strategy_zh.jpg';
 import { ResetConfigsButton } from '../components/ResetConfigsButton';
+import { useMediaQuery } from 'usehooks-ts';

 export type ApiParameters = {
  apiPort: number
@@ -30,7 +31,7 @@ export type ApiParameters = {
  frequencyPenalty: number;
 }

-export type Device = 'CPU' | 'CUDA' | 'MPS' | 'Custom';
+export type Device = 'CPU' | 'CUDA' | 'CUDA-Beta' | 'WebGPU' | 'MPS' | 'Custom';
 export type Precision = 'fp16' | 'int8' | 'fp32';

 export type ModelParameters = {
@@ -56,6 +57,7 @@ export const Configs: FC = observer(() => {
  const [selectedIndex, setSelectedIndex] = React.useState(commonStore.currentModelConfigIndex);
  const [selectedConfig, setSelectedConfig] = React.useState(commonStore.modelConfigs[selectedIndex]);
  const [displayStrategyImg, setDisplayStrategyImg] = React.useState(false);
+  const mq = useMediaQuery('(min-width: 640px)');
  const navigate = useNavigate();
  const port = selectedConfig.apiParameters.apiPort;

@@ -128,7 +130,8 @@ export const Configs: FC = observer(() => {
            setSelectedIndex(0);
            setSelectedConfig(commonStore.modelConfigs[0]);
          }} />
-          <ToolTipButton desc={t('Save Config')} icon={<Save20Regular />} onClick={onClickSave} />
+          <ToolTipButton desc={mq ? '' : t('Save Config')} icon={<Save20Regular />} text={mq ? t('Save Config') : ''}
+            onClick={onClickSave} />
        </div>
        <div className="flex items-center gap-4">
          <Label>{t('Config Name')}</Label>
@@ -237,40 +240,84 @@ export const Configs: FC = observer(() => {
                    }} />
                  </div>
                } />
-                <ToolTipButton text={t('Convert')}
-                  desc={t('Convert model with these configs. Using a converted model will greatly improve the loading speed, but model parameters of the converted model cannot be modified.')}
-                  onClick={async () => {
-                    if (commonStore.platform == 'darwin') {
-                      toast(t('MacOS is not yet supported for performing this operation, please do it manually.'), { type: 'info' });
-                      return;
-                    } else if (commonStore.platform == 'linux') {
-                      toast(t('Linux is not yet supported for performing this operation, please do it manually.'), { type: 'info' });
-                      return;
-                    }
-
-                    const modelPath = `${commonStore.settings.customModelsPath}/${selectedConfig.modelParameters.modelName}`;
-                    if (await FileExists(modelPath)) {
-                      const strategy = getStrategy(selectedConfig);
-                      const newModelPath = modelPath + '-' + strategy.replace(/[:> *+]/g, '-');
-                      toast(t('Start Converting'), { autoClose: 1000, type: 'info' });
-                      ConvertModel(commonStore.settings.customPythonPath, modelPath, strategy, newModelPath).then(async () => {
-                        if (!await FileExists(newModelPath + '.pth')) {
-                          toast(t('Convert Failed') + ' - ' + await GetPyError(), { type: 'error' });
-                        } else {
-                          toast(`${t('Convert Success')} - ${newModelPath}`, { type: 'success' });
+                {
+                  selectedConfig.modelParameters.device !== 'WebGPU' ?
+                    <ToolTipButton text={t('Convert')}
+                      desc={t('Convert model with these configs. Using a converted model will greatly improve the loading speed, but model parameters of the converted model cannot be modified.')}
+                      onClick={async () => {
+                        if (commonStore.platform === 'darwin') {
+                          toast(t('MacOS is not yet supported for performing this operation, please do it manually.') + ' (backend-python/convert_model.py)', { type: 'info' });
+                          return;
+                        } else if (commonStore.platform === 'linux') {
+                          toast(t('Linux is not yet supported for performing this operation, please do it manually.') + ' (backend-python/convert_model.py)', { type: 'info' });
+                          return;
                        }
-                      }).catch(e => {
-                        const errMsg = e.message || e;
-                        if (errMsg.includes('path contains space'))
-                          toast(`${t('Convert Failed')} - ${t('File Path Cannot Contain Space')}`, { type: 'error' });
-                        else
-                          toast(`${t('Convert Failed')} - ${e.message || e}`, { type: 'error' });
-                      });
-                      setTimeout(WindowShow, 1000);
-                    } else {
-                      toast(`${t('Model Not Found')} - ${modelPath}`, { type: 'error' });
-                    }
-                  }} />
+
+                        const ok = await checkDependencies(navigate);
+                        if (!ok)
+                          return;
+
+                        const modelPath = `${commonStore.settings.customModelsPath}/${selectedConfig.modelParameters.modelName}`;
+                        if (await FileExists(modelPath)) {
+                          const strategy = getStrategy(selectedConfig);
+                          const newModelPath = modelPath + '-' + strategy.replace(/[:> *+]/g, '-');
+                          toast(t('Start Converting'), { autoClose: 1000, type: 'info' });
+                          ConvertModel(commonStore.settings.customPythonPath, modelPath, strategy, newModelPath).then(async () => {
+                            if (!await FileExists(newModelPath + '.pth')) {
+                              toast(t('Convert Failed') + ' - ' + await GetPyError(), { type: 'error' });
+                            } else {
+                              toast(`${t('Convert Success')} - ${newModelPath}`, { type: 'success' });
+                            }
+                          }).catch(e => {
+                            const errMsg = e.message || e;
+                            if (errMsg.includes('path contains space'))
+                              toast(`${t('Convert Failed')} - ${t('File Path Cannot Contain Space')}`, { type: 'error' });
+                            else
+                              toast(`${t('Convert Failed')} - ${e.message || e}`, { type: 'error' });
+                          });
+                          setTimeout(WindowShow, 1000);
+                        } else {
+                          toast(`${t('Model Not Found')} - ${modelPath}`, { type: 'error' });
+                        }
+                      }} /> :
+                    <ToolTipButton text={t('Convert To Safe Tensors Format')}
+                      desc=""
+                      onClick={async () => {
+                        if (commonStore.platform === 'darwin') {
+                          toast(t('MacOS is not yet supported for performing this operation, please do it manually.') + ' (backend-python/convert_safetensors.py)', { type: 'info' });
+                          return;
+                        } else if (commonStore.platform === 'linux') {
+                          toast(t('Linux is not yet supported for performing this operation, please do it manually.') + ' (backend-python/convert_safetensors.py)', { type: 'info' });
+                          return;
+                        }
+
+                        const ok = await checkDependencies(navigate);
+                        if (!ok)
+                          return;
+
+                        const modelPath = `${commonStore.settings.customModelsPath}/${selectedConfig.modelParameters.modelName}`;
+                        if (await FileExists(modelPath)) {
+                          toast(t('Start Converting'), { autoClose: 1000, type: 'info' });
+                          const newModelPath = modelPath.replace(/\.pth$/, '.st');
+                          ConvertSafetensors(commonStore.settings.customPythonPath, modelPath, newModelPath).then(async () => {
+                            if (!await FileExists(newModelPath)) {
+                              toast(t('Convert Failed') + ' - ' + await GetPyError(), { type: 'error' });
+                            } else {
+                              toast(`${t('Convert Success')} - ${newModelPath}`, { type: 'success' });
+                            }
+                          }).catch(e => {
+                            const errMsg = e.message || e;
+                            if (errMsg.includes('path contains space'))
+                              toast(`${t('Convert Failed')} - ${t('File Path Cannot Contain Space')}`, { type: 'error' });
+                            else
+                              toast(`${t('Convert Failed')} - ${e.message || e}`, { type: 'error' });
+                          });
+                          setTimeout(WindowShow, 1000);
+                        } else {
+                          toast(`${t('Model Not Found')} - ${modelPath}`, { type: 'error' });
+                        }
+                      }} />
+                }
                <Labeled label={t('Strategy')} content={
                  <Dropdown style={{ minWidth: 0 }} className="grow" value={t(selectedConfig.modelParameters.device)!}
                    selectedOptions={[selectedConfig.modelParameters.device]}
@@ -284,11 +331,13 @@ export const Configs: FC = observer(() => {
                    <Option value="CPU">CPU</Option>
                    {commonStore.platform === 'darwin' && <Option value="MPS">MPS</Option>}
                    <Option value="CUDA">CUDA</Option>
+                    <Option value="CUDA-Beta">{t('CUDA (Beta, Faster)')!}</Option>
+                    <Option value="WebGPU">WebGPU</Option>
                    <Option value="Custom">{t('Custom')!}</Option>
                  </Dropdown>
                } />
                {
-                  selectedConfig.modelParameters.device != 'Custom' && <Labeled label={t('Precision')}
+                  selectedConfig.modelParameters.device !== 'Custom' && <Labeled label={t('Precision')}
                    desc={t('int8 uses less VRAM, but has slightly lower quality. fp16 has higher quality, and fp32 has the best quality.')}
                    content={
                      <Dropdown style={{ minWidth: 0 }} className="grow"
@@ -303,17 +352,17 @@ export const Configs: FC = observer(() => {
                        }}>
                        <Option>fp16</Option>
                        <Option>int8</Option>
-                        <Option>fp32</Option>
+                        {selectedConfig.modelParameters.device !== 'WebGPU' && <Option>fp32</Option>}
                      </Dropdown>
                    } />
                }
                {
-                  selectedConfig.modelParameters.device == 'CUDA' &&
+                  selectedConfig.modelParameters.device.includes('CUDA') &&
                  <Labeled label={t('Current Strategy')}
                    content={<Text> {getStrategy(selectedConfig)} </Text>} />
                }
                {
-                  selectedConfig.modelParameters.device == 'CUDA' &&
+                  selectedConfig.modelParameters.device.includes('CUDA') &&
                  <Labeled label={t('Stored Layers')}
                    desc={t('Number of the neural network layers loaded into VRAM, the more you load, the faster the speed, but it consumes more VRAM. (If your VRAM is not enough, it will fail to load)')}
                    content={
@@ -326,9 +375,7 @@ export const Configs: FC = observer(() => {
                        }} />
                    } />
                }
-                {
-                  selectedConfig.modelParameters.device == 'CUDA' && <div />
-                }
+                {selectedConfig.modelParameters.device.includes('CUDA') && <div />}
                {
                  displayStrategyImg &&
                  <img style={{ width: '80vh', height: 'auto', zIndex: 100 }}
@@ -336,13 +383,13 @@ export const Configs: FC = observer(() => {
                    src={commonStore.settings.language === 'zh' ? strategyZhImg : strategyImg} />
                }
                {
-                  selectedConfig.modelParameters.device == 'Custom' &&
+                  selectedConfig.modelParameters.device === 'Custom' &&
                  <Labeled label="Strategy"
                    onMouseEnter={() => setDisplayStrategyImg(true)}
                    onMouseLeave={() => setDisplayStrategyImg(false)}
                    content={
                      <Input className="grow"
-                        placeholder={commonStore.platform != 'darwin' ? 'cuda:0 fp16 *20 -> cuda:1 fp16' : 'mps fp32'}
+                        placeholder={commonStore.platform !== 'darwin' ? 'cuda:0 fp16 *20 -> cuda:1 fp16' : 'mps fp32'}
                        value={selectedConfig.modelParameters.customStrategy}
                        onChange={(e, data) => {
                          setSelectedConfigModelParams({
@@ -351,9 +398,9 @@ export const Configs: FC = observer(() => {
                        }} />
                    } />
                }
-                {selectedConfig.modelParameters.device == 'Custom' && <div />}
+                {selectedConfig.modelParameters.device === 'Custom' && <div />}
                {
-                  selectedConfig.modelParameters.device != 'CPU' && selectedConfig.modelParameters.device != 'MPS' &&
+                  (selectedConfig.modelParameters.device.includes('CUDA') || selectedConfig.modelParameters.device === 'Custom') &&
                  <Labeled label={t('Use Custom CUDA kernel to Accelerate')}
                    desc={t('Enabling this option can greatly improve inference speed and save some VRAM, but there may be compatibility issues. If it fails to start, please turn off this option.')}
                    content={
--- a/frontend/src/pages/Settings.tsx
+++ b/frontend/src/pages/Settings.tsx
@@ -126,7 +126,7 @@ export const Settings: FC = observer(() => {
          } />
        }
        {
-          commonStore.settings.language === 'zh' && commonStore.platform != 'linux' &&
+          commonStore.settings.language === 'zh' && commonStore.platform !== 'linux' &&
          <Labeled label={t('Use Tsinghua Pip Mirrors')} flex spaceBetween content={
            <Switch checked={commonStore.settings.cnMirror}
              onChange={(e, data) => {
--- a/frontend/src/startup.ts
+++ b/frontend/src/startup.ts
@@ -70,7 +70,7 @@ async function initConfig() {
      configData.currentModelConfigIndex >= 0 && configData.currentModelConfigIndex < configData.modelConfigs.length)
      commonStore.setCurrentConfigIndex(configData.currentModelConfigIndex, false);
  }).catch(() => {
-    commonStore.setModelConfigs(commonStore.platform != 'darwin' ? defaultModelConfigs : defaultModelConfigsMac, true);
+    commonStore.setModelConfigs(commonStore.platform !== 'darwin' ? defaultModelConfigs : defaultModelConfigsMac, true);
  });
 }

--- a/frontend/src/stores/commonStore.ts
+++ b/frontend/src/stores/commonStore.ts
@@ -167,7 +167,7 @@ class CommonStore {
  createModelConfig = (config: ModelConfig = defaultModelConfigs[0], saveConfig: boolean = true) => {
    if (config.name === defaultModelConfigs[0].name) {
      // deep copy
-      config = JSON.parse(JSON.stringify(commonStore.platform != 'darwin' ? defaultModelConfigs[0] : defaultModelConfigsMac[0]));
+      config = JSON.parse(JSON.stringify(commonStore.platform !== 'darwin' ? defaultModelConfigs[0] : defaultModelConfigsMac[0]));
      config.name = new Date().toLocaleString();
    }
    this.modelConfigs.push(config);
--- a/frontend/src/utils/index.tsx
+++ b/frontend/src/utils/index.tsx
@@ -57,6 +57,8 @@ export async function refreshBuiltInModels(readCache: boolean = false) {
  return cache;
 }

+const modelSuffix = ['.pth', '.st', '.safetensors'];
+
 export async function refreshLocalModels(cache: {
  models: ModelSourceItem[]
 }, filter: boolean = true, initUnfinishedModels: boolean = false) {
@@ -65,7 +67,7 @@ export async function refreshLocalModels(cache: {

  await ListDirFiles(commonStore.settings.customModelsPath).then((data) => {
    cache.models.push(...data.flatMap(d => {
-      if (!d.isDir && d.name.endsWith('.pth'))
+      if (!d.isDir && modelSuffix.some((ext => d.name.endsWith(ext))))
        return [{
          name: d.name,
          size: d.size,
@@ -146,7 +148,7 @@ export async function refreshRemoteModels(cache: { models: ModelSourceItem[] })
  .catch(() => {
  });
  cache.models = cache.models.filter((model, index, self) => {
-    return model.name.endsWith('.pth')
+    return modelSuffix.some((ext => model.name.endsWith(ext)))
      && index === self.findIndex(
        m => m.name === model.name || (m.SHA256 && m.SHA256 === model.SHA256 && m.size === model.size));
  });
@@ -176,7 +178,11 @@ export const getStrategy = (modelConfig: ModelConfig | undefined = undefined) =>
      strategy += 'cpu ';
      strategy += params.precision === 'int8' ? 'fp32i8' : 'fp32';
      break;
+    case 'WebGPU':
+      strategy += params.precision === 'int8' ? 'fp16i8' : 'fp16';
+      break;
    case 'CUDA':
+    case 'CUDA-Beta':
      if (avoidOverflow)
        strategy = 'cuda fp32 *1 -> ';
      strategy += 'cuda ';
@@ -239,7 +245,7 @@ export function downloadProgramFiles() {
  manifest.programFiles.forEach(({ url, path }) => {
    if (path)
      ReadFileInfo(path).then(info => {
-        if (info.size == 0 && url)
+        if (info.size === 0 && url)
          AddToDownloadList(path, url.replace('@master', '@v' + manifest.version));
      }).catch(() => {
        if (url)
@@ -372,7 +378,7 @@ export const checkDependencies = async (navigate: NavigateFunction) => {
          });
        } else {
          toast(depErrorMsg, { type: 'info', position: 'bottom-left' });
-          if (commonStore.platform != 'linux')
+          if (commonStore.platform !== 'linux')
            toastWithButton(t('Python dependencies are incomplete, would you like to install them?'), t('Install'), () => {
              InstallPyDep(commonStore.settings.customPythonPath, commonStore.settings.cnMirror).catch((e) => {
                const errMsg = e.message || e;
--- a/frontend/wailsjs/go/backend_golang/App.d.ts
+++ b/frontend/wailsjs/go/backend_golang/App.d.ts
@@ -10,6 +10,8 @@ export function ConvertData(arg1:string,arg2:string,arg3:string,arg4:string):Pro

 export function ConvertModel(arg1:string,arg2:string,arg3:string,arg4:string):Promise<string>;

+export function ConvertSafetensors(arg1:string,arg2:string,arg3:string):Promise<string>;
+
 export function CopyFile(arg1:string,arg2:string):Promise<void>;

 export function DeleteFile(arg1:string):Promise<void>;
@@ -46,7 +48,9 @@ export function RestartApp():Promise<void>;

 export function SaveJson(arg1:string,arg2:any):Promise<void>;

-export function StartServer(arg1:string,arg2:number,arg3:string):Promise<string>;
+export function StartServer(arg1:string,arg2:number,arg3:string,arg4:boolean):Promise<string>;
+
+export function StartWebGPUServer(arg1:number,arg2:string):Promise<string>;

 export function UpdateApp(arg1:string):Promise<boolean>;

--- a/frontend/wailsjs/go/backend_golang/App.js
+++ b/frontend/wailsjs/go/backend_golang/App.js
@@ -18,6 +18,10 @@ export function ConvertModel(arg1, arg2, arg3, arg4) {
  return window['go']['backend_golang']['App']['ConvertModel'](arg1, arg2, arg3, arg4);
 }

+export function ConvertSafetensors(arg1, arg2, arg3) {
+  return window['go']['backend_golang']['App']['ConvertSafetensors'](arg1, arg2, arg3);
+}
+
 export function CopyFile(arg1, arg2) {
  return window['go']['backend_golang']['App']['CopyFile'](arg1, arg2);
 }
@@ -90,8 +94,12 @@ export function SaveJson(arg1, arg2) {
  return window['go']['backend_golang']['App']['SaveJson'](arg1, arg2);
 }

-export function StartServer(arg1, arg2, arg3) {
-  return window['go']['backend_golang']['App']['StartServer'](arg1, arg2, arg3);
+export function StartServer(arg1, arg2, arg3, arg4) {
+  return window['go']['backend_golang']['App']['StartServer'](arg1, arg2, arg3, arg4);
+}
+
+export function StartWebGPUServer(arg1, arg2) {
+  return window['go']['backend_golang']['App']['StartWebGPUServer'](arg1, arg2);
 }

 export function UpdateApp(arg1) {
--- a/main.go
+++ b/main.go
@@ -49,6 +49,9 @@ var cyacInfo embed.FS
 //go:embed backend-python
 var py embed.FS

+//go:embed backend-rust
+var webgpu embed.FS
+
 //go:embed finetune
 var finetune embed.FS

@@ -63,6 +66,7 @@ func main() {
 		backend.CopyEmbed(cyac)
 		backend.CopyEmbed(cyacInfo)
 		backend.CopyEmbed(py)
+		backend.CopyEmbed(webgpu)
 		backend.CopyEmbed(finetune)
 		backend.CopyEmbed(midi)
 		backend.CopyEmbed(midiAssets)
--- a/manifest.json
+++ b/manifest.json
@@ -1,5 +1,5 @@
 {
-  "version": "1.4.2",
+  "version": "1.4.3",
  "introduction": {
    "en": "RWKV is an open-source, commercially usable large language model with high flexibility and great potential for development.\n### About This Tool\nThis tool aims to lower the barrier of entry for using large language models, making it accessible to everyone. It provides fully automated dependency and model management. You simply need to click and run, following the instructions, to deploy a local large language model. The tool itself is very compact and only requires a single executable file for one-click deployment.\nAdditionally, this tool offers an interface that is fully compatible with the OpenAI API. This means you can use any ChatGPT client as a client for RWKV, enabling capability expansion beyond just chat functionality.\n### Preset Configuration Rules at the Bottom\nThis tool comes with a series of preset configurations to reduce complexity. The naming rules for each configuration represent the following in order: device - required VRAM/memory - model size - model language.\nFor example, \"GPU-8G-3B-EN\" indicates that this configuration is for a graphics card with 8GB of VRAM, a model size of 3 billion parameters, and it uses an English language model.\nLarger model sizes have higher performance and VRAM requirements. Among configurations with the same model size, those with higher VRAM usage will have faster runtime.\nFor example, if you have 12GB of VRAM but running the \"GPU-12G-7B-EN\" configuration is slow, you can downgrade to \"GPU-8G-3B-EN\" for a significant speed improvement.\n### About RWKV\nRWKV is an RNN with Transformer-level LLM performance, which can also be directly trained like a GPT transformer (parallelizable). And it's 100% attention-free. You only need the hidden state at position t to compute the state at position t+1. You can use the \"GPT\" mode to quickly compute the hidden state for the \"RNN\" mode.<br/>So it's combining the best of RNN and transformer - great performance, fast inference, saves VRAM, fast training, \"infinite\" ctx_len, and free sentence embedding (using the final hidden state).",
    "zh": "RWKV是一个开源且允许商用的大语言模型，灵活性很高且极具发展潜力。\n### 关于本工具\n本工具旨在降低大语言模型的使用门槛，做到人人可用，本工具提供了全自动化的依赖和模型管理，你只需要直接点击运行，跟随引导，即可完成本地大语言模型的部署，工具本身体积极小，只需要一个exe即可完成一键部署。\n此外，本工具提供了与OpenAI API完全兼容的接口，这意味着你可以把任意ChatGPT客户端用作RWKV的客户端，实现能力拓展，而不局限于聊天。\n### 底部的预设配置规则\n本工具内置了一系列预设配置，以降低使用难度，每个配置名的规则，依次代表着：设备-所需显存/内存-模型规模-模型语言。\n例如，GPU-8G-3B-CN，表示该配置用于显卡，需要8G显存，模型规模为30亿参数，使用的是中文模型。\n模型规模越大，性能要求越高，显存要求也越高，而同样模型规模的配置中，显存占用越高的，运行速度越快。\n例如当你有12G显存，但运行GPU-12G-7B-CN配置速度比较慢，可降级成GPU-8G-3B-CN，将会大幅提速。\n### 关于RWKV\nRWKV是具有Transformer级别LLM性能的RNN，也可以像GPT Transformer一样直接进行训练（可并行化）。而且它是100% attention-free的。你只需在位置t处获得隐藏状态即可计算位置t + 1处的状态。你可以使用“GPT”模式快速计算用于“RNN”模式的隐藏状态。\n因此，它将RNN和Transformer的优点结合起来 - 高性能、快速推理、节省显存、快速训练、“无限”上下文长度以及免费的语句嵌入（使用最终隐藏状态）。"
--- a/vendor.yml
+++ b/vendor.yml
@@ -2,6 +2,7 @@
 - ^backend-python/wkv_cuda_utils/
 - ^backend-python/get-pip\.py
 - ^backend-python/convert_model\.py
+- ^backend-python/convert_safetensors\.py
 - ^backend-python/utils/midi\.py
 - ^build/
 - ^finetune/lora/
Author	SHA1	Message	Date
josc146	c9e4ae7fa1	release v1.4.4	2023-08-16 23:33:22 +08:00
josc146	79a97b2bc4	webgpu release support	2023-08-16 23:31:04 +08:00
josc146	ef53951a16	webgpu support	2023-08-16 23:07:58 +08:00
josc146	74f1a1c033	chore	2023-08-16 21:11:58 +08:00
josc146	ce986cfc6d	chore	2023-08-16 12:50:22 +08:00
josc146	61cea2a784	add misc API (`/models` and `/dashboard/billing/credit_grants`)	2023-08-14 23:37:55 +08:00
josc146	8a13bd3c1e	add rwkv-cuda-beta support (faster)	2023-08-14 22:07:15 +08:00
josc146	da68926e9c	chore (AddStateBody class)	2023-08-13 21:27:29 +08:00
josc146	e0b7453883	allow multiple systems	2023-08-04 22:27:55 +08:00
josc146	91e2828a95	allow completions input to be null	2023-08-04 22:22:59 +08:00
github-actions[bot]	bcf6409536	release v1.4.3	2023-07-31 14:51:01 +00:00