DeepSeek-V4-Flash 量化部署失败

Holferzhang
2026年7月13日
chat_bubble_outline 4
  • link
    Members 2 posts
    2026年7月13日 01:20
    link

    我下载了 modelscope.cn/models/metax-tech/DeepSeek-V4-Flash-FlexSMQ-AWQ-W8A8 这个量化模型，并且服务器上安装了最新Maca SDK，也下载了最新的vllm-metax镜像，配置 如下：
    [root@tcs-10-234-40-68 DeepSeek-V4-Flash-FlexSMQ-AWQ-W8A8]# mx-smi
    mx-smi version: 2.3.4

    =================== MetaX System Management Interface Log ===================
    Timestamp : Mon Jul 13 01:26:52 2026

    Attached GPUs : 8
    +---------------------------------------------------------------------------------+
    | MX-SMI 2.3.4 Kernel Mode Driver Version: 3.9.10 |
    | MACA Version: 3.8.0.23 BIOS Version: 1.35.2.0 |
    |------------------+-----------------+---------------------+----------------------|
    | Board Name | GPU Persist-M | Bus-id | GPU-Util sGPU-M |
    | Pwr:Usage/Cap | Temp Perf | Memory-Usage | GPU-State |
    |==================+=================+=====================+======================|
    | 0 MetaX C500 | 0 Off | 0000:1b:00.0 | 0% Disabled |
    | 37W / 350W | 38C P0 | 860/65536 MiB | Available |
    +------------------+-----------------+---------------------+----------------------+
    | 1 MetaX C500 | 1 Off | 0000:1c:00.0 | 0% Disabled |
    | 39W / 350W | 39C P0 | 860/65536 MiB | Available |
    +------------------+-----------------+---------------------+----------------------+
    | 2 MetaX C500 | 2 Off | 0000:21:00.0 | 0% Disabled |
    | 41W / 350W | 39C P0 | 860/65536 MiB | Available |
    +------------------+-----------------+---------------------+----------------------+
    | 3 MetaX C500 | 3 Off | 0000:24:00.0 | 0% Disabled |
    | 42W / 350W | 38C P0 | 860/65536 MiB | Available |
    +------------------+-----------------+---------------------+----------------------+
    | 4 MetaX C500 | 4 Off | 0000:2e:00.0 | 0% Disabled |
    | 40W / 350W | 38C P0 | 860/65536 MiB | Available |
    +------------------+-----------------+---------------------+----------------------+
    | 5 MetaX C500 | 5 Off | 0000:34:00.0 | 0% Disabled |
    | 40W / 350W | 38C P0 | 860/65536 MiB | Available |
    +------------------+-----------------+---------------------+----------------------+
    | 6 MetaX C500 | 6 Off | 0000:37:00.0 | 0% Disabled |
    | 41W / 350W | 40C P0 | 860/65536 MiB | Available |
    +------------------+-----------------+---------------------+----------------------+
    | 7 MetaX C500 | 7 Off | 0000:38:00.0 | 0% Disabled |
    | 47W / 350W | 43C P0 | 860/65536 MiB | Available |
    +------------------+-----------------+---------------------+----------------------+

    +---------------------------------------------------------------------------------+
    | Process: |
    | GPU PID Process Name GPU Memory |
    | Usage(MiB) |
    |=================================================================================|
    | no process found |
    +---------------------------------------------------------------------------------+

    End of Log

    [root@tcs-10-234-40-68 DeepSeek-V4-Flash-FlexSMQ-AWQ-W8A8]# docker ps | grep mx
    18c870cbacae cr.metax-tech.com/public-ai-release/maca/vllm-metax:0.21.0-maca.ai3.7.1.106-torch2.8-py310-ubuntu22.04-amd64 "/bin/bash" 2 hours ago Up 2 hours 0.0.0.0:15570->5570/tcp, 0.0.0.0:8002->8000/tcp mx_vllm_deepseek_v4_flash

    vllm的启动脚本如下：
    [root@tcs-10-234-40-68 DeepSeek-V4-Flash-FlexSMQ-AWQ-W8A8]# cat /data1/vllm/userdata/holfer/python_scripts/run_deepseek_v4_flash.sh

    !/bin/bash

    1. 核心关键：强制关闭 vLLM 的 V1 架构，回退到稳定的 V0 架构（避开 V1 的强制 TileLang 算子流）

    export VLLM_USE_V1=0

    2. 强行关闭任何潜在的 Torch Compile 编译行为

    export ENABLE_DEVICE_COMPILE=0
    export TILELANG_ENABLE_CACHE=0
    export VLLM_TORCH_COMPILE_LEVEL=0

    3. 保持沐曦原有的加速环境变量

    export MACA_SMALL_PAGESIZE_ENABLE=1
    export MACA_VLLM_ENABLE_MCTLASS_FUSED_MOE=1
    export MACA_VLLM_ENABLE_MCTLASS_PYTHON_API=1

    4. 启动服务（保持你的原始参数，确保 --enforce-eager 在线）

    VLLM_USE_MODELSCOPE=false vllm serve /holfer/model/metax-tech/DeepSeek-V4-Flash-FlexSMQ-AWQ-W8A8 \
    --port 8000 \
    --trust-remote-code \
    --kv-cache-dtype bfloat16 \
    --block-size 256 \
    --gpu-memory-utilization 0.80 \
    --enforce-eager \
    --served-model-name deepseek-v4-flash \
    --tensor-parallel-size 8 \
    --dtype bfloat16 \
    --tokenizer-mode deepseek_v4 \
    --tool-call-parser deepseek_v4 \
    --enable-auto-tool-choice \
    --reasoning-parser deepseek_v4 \
    --max-num-seq 32 \
    --speculative_config '{"method": "mtp", "num_speculative_tokens": 1}'
    [root@tcs-10-234-40-68 DeepSeek-V4-Flash-FlexSMQ-AWQ-W8A8]#

    一直报错： (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] Compilation error:
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] mxcc: warning: argument unused during compilation: '-lruntime_cu' [-Wunused-command-line-argument]
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] mxcc: warning: argument unused during compilation: '-lToolsExt_cu' [-Wunused-command-line-argument]
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] In file included from /tmp/tmpdjspkits/tvm_kernels.cu:3:
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] In file included from /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/copy.h:3:
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] In file included from /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/common.h:9:
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:86:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 86 | : "=h"(ret)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:91:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 91 | : "=h"(ret)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:96:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 96 | : "=h"(ret)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:106:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 106 | : "=h"(ret)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:111:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 111 | : "=h"(ret)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:116:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 116 | : "=h"(ret)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:130:11: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 130 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:130:24: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 130 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:136:11: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 136 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:136:24: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 136 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:142:11: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 142 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:142:24: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 142 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:155:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 155 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:155:33: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 155 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:160:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 160 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:160:33: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 160 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:165:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 165 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:165:33: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 165 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:176:20: error: invalid output constraint '=f' in asm
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 176 | : "=f"(ret_x), "=f"(ret_y)
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] | ^
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] fatal error: too many errors emitted, stopping now [-ferror-limit=]
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] 20 errors generated when compiling for xcore1000.
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140]
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] Command: /opt/maca/tools/cu-bridge/bin/nvcc --cubin -O3 -lineinfo -arch=sm_80 -std=c++17 -I/opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src -I/opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/cutlass/include --ptxas-options=--register-usage-level=10 -o /tmp/tmpdjspkits/tvm_kernels.cubin /tmp/tmpdjspkits/tvm_kernels.cu
    (EngineCore pid=34080) ERROR 07-13 01:20:47 [core.py:1140] ', please check the stack trace above for the root cause
    (Worker_TP0 pid=34220) WARNING 07-13 01:20:47 [multiproc_executor.py:884] WorkerProc was terminated
    (Worker_TP7 pid=34257) WARNING 07-13 01:20:47 [multiproc_executor.py:884] WorkerProc was terminated
    (Worker_TP3 pid=34226) WARNING 07-13 01:20:47 [multiproc_executor.py:884] WorkerProc was terminated
    (Worker_TP2 pid=34224) WARNING 07-13 01:20:47 [multiproc_executor.py:884] WorkerProc was terminated
    (Worker_TP4 pid=34230) WARNING 07-13 01:20:47 [multiproc_executor.py:884] WorkerProc was terminated
    (Worker_TP5 pid=34239) WARNING 07-13 01:20:47 [multiproc_executor.py:884] WorkerProc was terminated
    (EngineCore pid=34080) ERROR 07-13 01:20:56 [multiproc_executor.py:283] Worker proc VllmWorker-5 died unexpectedly, shutting down executor.
    (EngineCore pid=34080) Process EngineCore:
    (EngineCore pid=34080) Traceback (most recent call last):
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/multiprocessing/process.py", line 314, in _bootstrap
    (EngineCore pid=34080) self.run()
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/multiprocessing/process.py", line 108, in run
    (EngineCore pid=34080) self._target(self._args, self._kwargs)
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/core.py", line 1144, in run_engine_core
    (EngineCore pid=34080) raise e
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/core.py", line 1114, in run_engine_core
    (EngineCore pid=34080) engine_core = EngineCoreProc(    args, engine_index=dp_rank, kwargs)
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/tracing/otel.py", line 178, in sync_wrapper
    (EngineCore pid=34080) return func(*args,     kwargs)
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/core.py", line 880, in init
    (EngineCore pid=34080) super().init(
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/core.py", line 128, in init
    (EngineCore pid=34080) kv_cache_config = self._initialize_kv_caches(vllm_config)
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/tracing/otel.py", line 178, in sync_wrapper
    (EngineCore pid=34080) return func(args, kwargs)
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/core.py", line 250, in _initialize_kv_caches
    (EngineCore pid=34080) available_gpu_memory = self.model_executor.determine_available_memory()
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/executor/abstract.py", line 147, in determine_available_memory
    (EngineCore pid=34080) return self.collective_rpc("determine_available_memory")
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/executor/multiproc_executor.py", line 403, in collective_rpc
    (EngineCore pid=34080) return future if non_block else future.result()
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/executor/multiproc_executor.py", line 90, in result
    (EngineCore pid=34080) return super().result()
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/concurrent/futures/_base.py", line 451, in result
    (EngineCore pid=34080) return self.get_result()
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/concurrent/futures/_base.py", line 403, in __get_result
    (EngineCore pid=34080) raise self._exception
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/executor/multiproc_executor.py", line 94, in _wait_for_response
    (EngineCore pid=34080) response = self.aggregate(self.get_response())
    (EngineCore pid=34080) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/executor/multiproc_executor.py", line 390, in get_response
    (EngineCore pid=34080) raise RuntimeError(
    (EngineCore pid=34080) RuntimeError: Worker failed with error '#include <math_constants.h>
    (EngineCore pid=34080) #include <tl_templates/cuda/gemm.h>
    (EngineCore pid=34080) #include <tl_templates/cuda/copy.h>
    (EngineCore pid=34080) #include <tl_templates/cuda/reduce.h>
    (EngineCore pid=34080) #include <tl_templates/cuda/ldsm.h>
    (EngineCore pid=34080) #include <tl_templates/cuda/threadblock_swizzle.h>
    (EngineCore pid=34080) #include <tl_templates/cuda/debug.h>
    (EngineCore pid=34080) #ifdef ENABLE_BF16
    (EngineCore pid=34080) #include <tl_templates/cuda/cuda_bf16_fallbacks.cuh>
    (EngineCore pid=34080) #endif
    (EngineCore pid=34080)
    (EngineCore pid=34080) extern "C" __global     void mhc_pre_big_fuse_kernel(float     restrict comb_mix, const float restrict gemm_out_mul, const float restrict gemm_out_sqrsum, bfloat16_t restrict layer_input, const float restrict mhc_base, const float restrict mhc_scale, float restrict post_mix, const bfloat16_t restrict residual, int num_tokens);
    (EngineCore pid=34080) extern "C" global void launch_bounds(128, 1) mhc_pre_big_fuse_kernel(float     restrict comb_mix, const float restrict gemm_out_mul, const float restrict gemm_out_sqrsum, bfloat16_t restrict layer_input, const float restrict mhc_base, const float restrict mhc_scale, float restrict post_mix, const bfloat16_t restrict residual, int num_tokens) {
    (EngineCore pid=34080) extern shared align(1024) uchar buf_dyn_shmem[];
    (EngineCore pid=34080) float mixes[1];
    (EngineCore pid=34080) float rms[1];
    (EngineCore pid=34080) float cm[1];
    (EngineCore pid=34080) float row_max[1];
    (EngineCore pid=34080) float row_sum[1];
    (EngineCore pid=34080) float col_sum[1];
    (EngineCore pid=34080) float xl[32];
    (EngineCore pid=34080) float ol[8];
    (EngineCore pid=34080) bfloat16_t xs_local_cast[8];
    (EngineCore pid=34080) bfloat16_t layer_input_local_cast_1[8];
    (EngineCore pid=34080) bfloat16_t xs_local_cast_2[8];
    (EngineCore pid=34080) bfloat16_t layer_input_local_cast_3[8];
    (EngineCore pid=34080) bfloat16_t xs_local_cast_4[8];
    (EngineCore pid=34080) bfloat16_t layer_input_local_cast_5[8];
    (EngineCore pid=34080) if (((int)threadIdx.x) < 64) {
    (EngineCore pid=34080) if (((int)threadIdx.x) < 48) {
    (EngineCore pid=34080) mixes[0] = 0x0p+0f/    0.000000e+00/;
    (EngineCore pid=34080) }
    (EngineCore pid=34080) rms[0] = 0x0p+0f/    0.000000e+00/;
    (EngineCore pid=34080) for (int i_split = 0; i_split < 16; ++i_split) {
    (EngineCore pid=34080) rms[0] = (rms[0] + gemm_out_sqrsum[((((int64_t)i_split) * ((int64_t)num_tokens)) + ((int64_t)((int)blockIdx.x)))]);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) rms[0] = rsqrtf(((rms[0] / 0x1p+14f/    1.638400e+04/) + 0x1.0c6f7a0b5ed8dp-20f/1.000000e-06/));
    (EngineCore pid=34080) if (((int)threadIdx.x) < 48) {
    (EngineCore pid=34080) mixes[0] = 0x0p+0f/    0.000000e+00/;
    (EngineCore pid=34080) for (int i_split_1 = 0; i_split_1 < 16; ++i_split_1) {
    (EngineCore pid=34080) mixes[0] = (mixes[0] + gemm_out_mul[(((((int64_t)((int)blockIdx.x)) * (int64_t)24) + ((((int64_t)i_split_1) * ((int64_t)num_tokens)) * (int64_t)24)) + (((int64_t)((int)threadIdx.x)) % (int64_t)24))]);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) mixes[0] = (mixes[0] * rms[0]);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) if (((((int)threadIdx.x) / 24) == 0) && (((int)threadIdx.x) < 48)) {
    (EngineCore pid=34080) ((float    )buf_dyn_shmem)[(((int)threadIdx.x) % 24)] = mixes[0];
    (EngineCore pid=34080) }
    (EngineCore pid=34080) }
    (EngineCore pid=34080) syncthreads();
    (EngineCore pid=34080) if (((int)threadIdx.x) < 64) {
    (EngineCore pid=34080) if (((int)threadIdx.x) < 4) {
    (EngineCore pid=34080) post_mix[((((int64_t)((int)blockIdx.x)) * (int64_t)4) + ((int64_t)((int)threadIdx.x)))] = ((0x1p+0f/1.000000e+00/ / (0x1p+0f/1.000000e+00/ + expf((0x0p+0f/0.000000e+00/ - ((((float)buf_dyn_shmem)[(((int)threadIdx.x) + 4)] * mhc_scale[1]) + mhc_base[(((int)threadIdx.x) + 4)]))))) * 0x1p+1f/2.000000e+00/);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) cm[0] = ((((float    )buf_dyn_shmem)[((((int)threadIdx.x) & 15) + 8)] * mhc_scale[2]) + mhc_base[((((int)threadIdx.x) & 15) + 8)]);
    (EngineCore pid=34080) row_max[0] = -CUDART_INF_F;
    (EngineCore pid=34080) row_max[0] = max(row_max[0], cm[0]);
    (EngineCore pid=34080) row_max[0] = tl::AllReduce<tl::MaxOp, 4, 1, 0>::run(row_max[0]);
    (EngineCore pid=34080) cm[0] = expf((cm[0] - row_max[0]));
    (EngineCore pid=34080) row_sum[0] = 0x0p+0f/0.000000e+00/;
    (EngineCore pid=34080) row_sum[0] = (row_sum[0] + cm[0]);
    (EngineCore pid=34080) row_sum[0] = tl::AllReduce<tl::SumOp, 4, 1, 0>::run(row_sum[0]);
    (EngineCore pid=34080) cm[0] = ((cm[0] / row_sum[0]) + 0x1.0c6f7a0b5ed8dp-20f/1.000000e-06/);
    (EngineCore pid=34080) col_sum[0] = 0x0p+0f/0.000000e+00/;
    (EngineCore pid=34080) col_sum[0] = (col_sum[0] + cm[0]);
    (EngineCore pid=34080) col_sum[0] = tl::AllReduce<tl::SumOp, 16, 4, 0>::run(col_sum[0]);
    (EngineCore pid=34080) cm[0] = (cm[0] / (col_sum[0] + 0x1.0c6f7a0b5ed8dp-20f/1.000000e-06/));
    (EngineCore pid=34080) for (int __1 = 0; __1 < 19; ++__1) {
    (EngineCore pid=34080) row_sum[0] = 0x0p+0f/0.000000e+00/;
    (EngineCore pid=34080) row_sum[0] = (row_sum[0] + cm[0]);
    (EngineCore pid=34080) row_sum[0] = tl::AllReduce<tl::SumOp, 4, 1, 0>::run(row_sum[0]);
    (EngineCore pid=34080) cm[0] = (cm[0] / (row_sum[0] + 0x1.0c6f7a0b5ed8dp-20f/1.000000e-06/));
    (EngineCore pid=34080) col_sum[0] = 0x0p+0f/0.000000e+00/;
    (EngineCore pid=34080) col_sum[0] = (col_sum[0] + cm[0]);
    (EngineCore pid=34080) col_sum[0] = tl::AllReduce<tl::SumOp, 16, 4, 0>::run(col_sum[0]);
    (EngineCore pid=34080) cm[0] = (cm[0] / (col_sum[0] + 0x1.0c6f7a0b5ed8dp-20f/1.000000e-06/));
    (EngineCore pid=34080) }
    (EngineCore pid=34080) if ((((int)threadIdx.x) >> 4) == 0) {
    (EngineCore pid=34080) comb_mix[((((int64_t)((int)blockIdx.x)) * (int64_t)16) + (((int64_t)((int)threadIdx.x)) & (int64_t)15))] = cm[0];
    (EngineCore pid=34080) }
    (EngineCore pid=34080) } else {
    (EngineCore pid=34080) if (((int)threadIdx.x) < 68) {
    (EngineCore pid=34080) ((float)buf_dyn_shmem)[(((int)threadIdx.x) + 2008)] = ((0x1p+0f/1.000000e+00/ / (0x1p+0f/1.000000e+00/ + expf((0x0p+0f/0.000000e+00/ - ((((float)buf_dyn_shmem)[(((int)threadIdx.x) - 64)] * mhc_scale[0]) + mhc_base[(((int)threadIdx.x) - 64)]))))) + 0x1.0c6f7a0b5ed8dp-20f/1.000000e-06/);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) tl::__sync_thread_partial<3, 64>();
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i = 0; i < 4; ++i) {
    (EngineCore pid=34080) tl::cp_async_gs<16>((&(((bfloat16_t)buf_dyn_shmem)[(((i * 512) + (((int)threadIdx.x) * 8)) - 464)])), (&(residual[(((((((int64_t)((int)blockIdx.x)) * (int64_t)16384) + ((((int64_t)((int)threadIdx.x)) >> (int64_t)6) * (int64_t)4096)) + (((int64_t)i) * (int64_t)4096)) + ((((int64_t)((int)threadIdx.x)) & (int64_t)63) * (int64_t)8)) - (int64_t)4096)])));
    (EngineCore pid=34080) }
    (EngineCore pid=34080) tl::cp_async_commit();
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_1 = 0; i_1 < 4; ++i_1) {
    (EngineCore pid=34080) tl::cp_async_gs<16>((&(((bfloat16_t    )buf_dyn_shmem)[(((i_1 * 512) + (((int)threadIdx.x) * 8)) + 1584)])), (&(residual[(((((((int64_t)((int)blockIdx.x)) * (int64_t)16384) + ((((int64_t)((int)threadIdx.x)) >> (int64_t)6) * (int64_t)4096)) + (((int64_t)i_1) * (int64_t)4096)) + ((((int64_t)((int)threadIdx.x)) & (int64_t)63) * (int64_t)8)) - (int64_t)3584)])));
    (EngineCore pid=34080) }
    (EngineCore pid=34080) tl::cp_async_commit();
    (EngineCore pid=34080) for (int i0_h = 0; i0_h < 6; ++i0_h) {
    (EngineCore pid=34080) tl::cp_async_wait<1>();
    (EngineCore pid=34080) tl::__sync_thread_partial<3, 64>();
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_2 = 0; i_2 < 4; ++i_2) {
    (EngineCore pid=34080) (uint4)(xs_local_cast + 0) = (uint4)(((bfloat16_t)buf_dyn_shmem) + (((((i0_h & 1) * 2048) + (i_2 * 512)) + (((int)threadIdx.x) * 8)) - 464));
    (EngineCore pid=34080) for (int vec = 0; vec < 2; ++vec) {
    (EngineCore pid=34080) float4 __2;
    (EngineCore pid=34080) uint2 v_ =     (uint2)(xs_local_cast + (vec * 4));
    (EngineCore pid=34080) ((float2    )(&__2))[0] = __bfloat1622float2((reinterpret_cast<__nv_bfloat162>(&v_))[0]);
    (EngineCore pid=34080) ((float2    )(&__2))[1] = __bfloat1622float2((reinterpret_cast<__nv_bfloat162>(&v_))[1]);
    (EngineCore pid=34080)     (float4)(xl + ((i_2 * 8) + (vec * 4))) = __2;
    (EngineCore pid=34080) }
    (EngineCore pid=34080) }
    (EngineCore pid=34080) tl::__sync_thread_partial<3, 64>();
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_3 = 0; i_3 < 4; ++i_3) {
    (EngineCore pid=34080) tl::cp_async_gs<16>((&(((bfloat16_t    )buf_dyn_shmem)[(((((i0_h & 1) * 2048) + (i_3 * 512)) + (((int)threadIdx.x) * 8)) - 464)])), (&(residual[((((((((int64_t)((int)blockIdx.x)) * (int64_t)16384) + ((((int64_t)((int)threadIdx.x)) >> (int64_t)6) * (int64_t)4096)) + (((int64_t)i_3) * (int64_t)4096)) + (((int64_t)i0_h) * (int64_t)512)) + ((((int64_t)((int)threadIdx.x)) & (int64_t)63) * (int64_t)8)) - (int64_t)3072)])));
    (EngineCore pid=34080) }
    (EngineCore pid=34080) tl::cp_async_commit();
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_4 = 0; i_4 < 2; ++i_4) {
    (EngineCore pid=34080) float broadcast_var = 0x0p+0f/0.000000e+00/;
    (EngineCore pid=34080) (float4)(ol + (i_4 * 4)) = make_float4(broadcast_var, broadcast_var, broadcast_var, broadcast_var);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) tl::__sync_thread_partial<3, 64>();
    (EngineCore pid=34080) for (int i_mhc = 0; i_mhc < 4; ++i_mhc) {
    (EngineCore pid=34080) float pre = ((float)buf_dyn_shmem)[(i_mhc + 2072)];
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_5 = 0; i_5 < 8; ++i_5) {
    (EngineCore pid=34080) ol[i_5] = (ol[i_5] + (pre * xl[((i_mhc * 8) + i_5)]));
    (EngineCore pid=34080) }
    (EngineCore pid=34080) }
    (EngineCore pid=34080) for (int i_6 = 0; i_6 < 2; ++i_6) {
    (EngineCore pid=34080) uint2 __3;
    (EngineCore pid=34080) float4 v__1 =     (float4)(ol + (i_6 * 4));
    (EngineCore pid=34080) (reinterpret_cast<__nv_bfloat162    >(&__3))[0] = __float22bfloat162_rn(((float2)(&v__1))[0]);
    (EngineCore pid=34080) (reinterpret_cast<__nv_bfloat162    >(&__3))[1] = __float22bfloat162_rn(((float2)(&v__1))[1]);
    (EngineCore pid=34080)     (uint2)(layer_input_local_cast_1 + (i_6 * 4)) = __3;
    (EngineCore pid=34080) }
    (EngineCore pid=34080)     (uint4)(layer_input + ((((((int64_t)((int)blockIdx.x)) * (int64_t)4096) + (((int64_t)i0_h) * (int64_t)512)) + (((int64_t)((int)threadIdx.x)) * (int64_t)8)) - (int64_t)512)) = (uint4)(layer_input_local_cast_1 + 0);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) tl::cp_async_wait<1>();
    (EngineCore pid=34080) tl::__sync_thread_partial<3, 64>();
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_7 = 0; i_7 < 4; ++i_7) {
    (EngineCore pid=34080)     (uint4)(xs_local_cast_2 + 0) = (uint4)(((bfloat16_t)buf_dyn_shmem) + (((i_7 * 512) + (((int)threadIdx.x) * 8)) - 464));
    (EngineCore pid=34080) for (int vec_1 = 0; vec_1 < 2; ++vec_1) {
    (EngineCore pid=34080) float4 __4;
    (EngineCore pid=34080) uint2 v__2 = (uint2)(xs_local_cast_2 + (vec_1 * 4));
    (EngineCore pid=34080) ((float2)(&__4))[0] = __bfloat1622float2((reinterpret_cast<__nv_bfloat162>(&v__2))[0]);
    (EngineCore pid=34080) ((float2)(&__4))[1] = __bfloat1622float2((reinterpret_cast<__nv_bfloat162>(&v__2))[1]);
    (EngineCore pid=34080) (float4)(xl + ((i_7 * 8) + (vec_1 * 4))) = __4;
    (EngineCore pid=34080) }
    (EngineCore pid=34080) }
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_8 = 0; i_8 < 2; ++i_8) {
    (EngineCore pid=34080) float broadcast_var_1 = 0x0p+0f/0.000000e+00/;
    (EngineCore pid=34080) (float4)(ol + (i_8 * 4)) = make_float4(broadcast_var_1, broadcast_var_1, broadcast_var_1, broadcast_var_1);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) for (int i_mhc_1 = 0; i_mhc_1 < 4; ++i_mhc_1) {
    (EngineCore pid=34080) float pre_1 = ((float)buf_dyn_shmem)[(i_mhc_1 + 2072)];
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_9 = 0; i_9 < 8; ++i_9) {
    (EngineCore pid=34080) ol[i_9] = (ol[i_9] + (pre_1 * xl[((i_mhc_1 * 8) + i_9)]));
    (EngineCore pid=34080) }
    (EngineCore pid=34080) }
    (EngineCore pid=34080) for (int i_10 = 0; i_10 < 2; ++i_10) {
    (EngineCore pid=34080) uint2 __5;
    (EngineCore pid=34080) float4 v__3 =     (float4)(ol + (i_10 * 4));
    (EngineCore pid=34080) (reinterpret_cast<__nv_bfloat162    >(&__5))[0] = __float22bfloat162_rn(((float2)(&v__3))[0]);
    (EngineCore pid=34080) (reinterpret_cast<__nv_bfloat162    >(&__5))[1] = __float22bfloat162_rn(((float2)(&v__3))[1]);
    (EngineCore pid=34080)     (uint2)(layer_input_local_cast_3 + (i_10 * 4)) = __5;
    (EngineCore pid=34080) }
    (EngineCore pid=34080)     (uint4)(layer_input + (((((int64_t)((int)blockIdx.x)) * (int64_t)4096) + (((int64_t)((int)threadIdx.x)) * (int64_t)8)) + (int64_t)2560)) = (uint4)(layer_input_local_cast_3 + 0);
    (EngineCore pid=34080) tl::cp_async_wait<0>();
    (EngineCore pid=34080) tl::__sync_thread_partial<3, 64>();
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_11 = 0; i_11 < 4; ++i_11) {
    (EngineCore pid=34080)     (uint4)(xs_local_cast_4 + 0) = (uint4)(((bfloat16_t)buf_dyn_shmem) + (((i_11 * 512) + (((int)threadIdx.x) * 8)) + 1584));
    (EngineCore pid=34080) for (int vec_2 = 0; vec_2 < 2; ++vec_2) {
    (EngineCore pid=34080) float4 __6;
    (EngineCore pid=34080) uint2 v__4 = (uint2)(xs_local_cast_4 + (vec_2 * 4));
    (EngineCore pid=34080) ((float2)(&__6))[0] = __bfloat1622float2((reinterpret_cast<__nv_bfloat162>(&v__4))[0]);
    (EngineCore pid=34080) ((float2)(&__6))[1] = __bfloat1622float2((reinterpret_cast<__nv_bfloat162>(&v__4))[1]);
    (EngineCore pid=34080) (float4)(xl + ((i_11 * 8) + (vec_2 * 4))) = __6;
    (EngineCore pid=34080) }
    (EngineCore pid=34080) }
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_12 = 0; i_12 < 2; ++i_12) {
    (EngineCore pid=34080) float broadcast_var_2 = 0x0p+0f/0.000000e+00/;
    (EngineCore pid=34080) (float4)(ol + (i_12 * 4)) = make_float4(broadcast_var_2, broadcast_var_2, broadcast_var_2, broadcast_var_2);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) for (int i_mhc_2 = 0; i_mhc_2 < 4; ++i_mhc_2) {
    (EngineCore pid=34080) float pre_2 = ((float)buf_dyn_shmem)[(i_mhc_2 + 2072)];
    (EngineCore pid=34080) #pragma unroll
    (EngineCore pid=34080) for (int i_13 = 0; i_13 < 8; ++i_13) {
    (EngineCore pid=34080) ol[i_13] = (ol[i_13] + (pre_2 * xl[((i_mhc_2 * 8) + i_13)]));
    (EngineCore pid=34080) }
    (EngineCore pid=34080) }
    (EngineCore pid=34080) for (int i_14 = 0; i_14 < 2; ++i_14) {
    (EngineCore pid=34080) uint2 __7;
    (EngineCore pid=34080) float4 v__5 =     (float4)(ol + (i_14 * 4));
    (EngineCore pid=34080) (reinterpret_cast<__nv_bfloat162    >(&__7))[0] = __float22bfloat162_rn(((float2)(&v__5))[0]);
    (EngineCore pid=34080) (reinterpret_cast<__nv_bfloat162    >(&__7))[1] = __float22bfloat162_rn(((float2)(&v__5))[1]);
    (EngineCore pid=34080)     (uint2)(layer_input_local_cast_5 + (i_14 * 4)) = __7;
    (EngineCore pid=34080) }
    (EngineCore pid=34080)     (uint4)(layer_input + (((((int64_t)((int)blockIdx.x)) * (int64_t)4096) + (((int64_t)((int)threadIdx.x)) * (int64_t)8)) + (int64_t)3072)) = (uint4)(layer_input_local_cast_5 + 0);
    (EngineCore pid=34080) }
    (EngineCore pid=34080) }
    (EngineCore pid=34080)
    (EngineCore pid=34080)
    (EngineCore pid=34080) Compilation error:
    (EngineCore pid=34080) mxcc: warning: argument unused during compilation: '-lruntime_cu' [-Wunused-command-line-argument]
    (EngineCore pid=34080) mxcc: warning: argument unused during compilation: '-lToolsExt_cu' [-Wunused-command-line-argument]
    (EngineCore pid=34080) In file included from /tmp/tmpdjspkits/tvm_kernels.cu:3:
    (EngineCore pid=34080) In file included from /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/copy.h:3:
    (EngineCore pid=34080) In file included from /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/common.h:9:
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:86:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 86 | : "=h"(ret)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:91:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 91 | : "=h"(ret)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:96:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 96 | : "=h"(ret)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:106:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 106 | : "=h"(ret)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:111:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 111 | : "=h"(ret)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:116:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 116 | : "=h"(ret)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:130:11: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 130 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:130:24: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 130 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:136:11: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 136 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:136:24: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 136 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:142:11: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 142 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:142:24: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 142 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:155:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 155 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:155:33: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 155 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:160:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 160 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:160:33: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 160 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:165:20: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 165 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:165:33: error: invalid output constraint '=h' in asm
    (EngineCore pid=34080) 165 | : "=h"(ret_x), "=h"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) /opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src/tl_templates/cuda/atomic.h:176:20: error: invalid output constraint '=f' in asm
    (EngineCore pid=34080) 176 | : "=f"(ret_x), "=f"(ret_y)
    (EngineCore pid=34080) | ^
    (EngineCore pid=34080) fatal error: too many errors emitted, stopping now [-ferror-limit=]
    (EngineCore pid=34080) 20 errors generated when compiling for xcore1000.
    (EngineCore pid=34080)
    (EngineCore pid=34080) Command: /opt/maca/tools/cu-bridge/bin/nvcc --cubin -O3 -lineinfo -arch=sm_80 -std=c++17 -I/opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/../src -I/opt/conda/lib/python3.10/site-packages/tilelang/3rdparty/cutlass/include --ptxas-options=--register-usage-level=10 -o /tmp/tmpdjspkits/tvm_kernels.cubin /tmp/tmpdjspkits/tvm_kernels.cu
    (EngineCore pid=34080) ', please check the stack trace above for the root cause
    (APIServer pid=33938) Traceback (most recent call last):
    (APIServer pid=33938) File "/opt/conda/bin/vllm", line 8, in <module>
    (APIServer pid=33938) sys.exit(main())
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/entrypoints/cli/main.py", line 92, in main
    (APIServer pid=33938) args.dispatch_function(args)
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/entrypoints/cli/serve.py", line 122, in cmd
    (APIServer pid=33938) uvloop.run(run_server(args))
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/uvloop/init.py", line 69, in run
    (APIServer pid=33938) return loop.run_until_complete(wrapper())
    (APIServer pid=33938) File "uvloop/loop.pyx", line 1518, in uvloop.loop.Loop.run_until_complete
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/uvloop/init.py", line 48, in wrapper
    (APIServer pid=33938) return await main
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/entrypoints/openai/api_server.py", line 693, in run_server
    (APIServer pid=33938) await run_server_worker(listen_address, sock, args,     uvicorn_kwargs)
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/entrypoints/openai/api_server.py", line 707, in run_server_worker
    (APIServer pid=33938) async with build_async_engine_client(
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/contextlib.py", line 199, in aenter
    (APIServer pid=33938) return await anext(self.gen)
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/entrypoints/openai/api_server.py", line 100, in build_async_engine_client
    (APIServer pid=33938) async with build_async_engine_client_from_engine_args(
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/contextlib.py", line 199, in aenter
    (APIServer pid=33938) return await anext(self.gen)
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/entrypoints/openai/api_server.py", line 136, in build_async_engine_client_from_engine_args
    (APIServer pid=33938) async_llm = AsyncLLM.from_vllm_config(
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/async_llm.py", line 217, in from_vllm_config
    (APIServer pid=33938) return cls(
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/async_llm.py", line 146, in init
    (APIServer pid=33938) self.engine_core = EngineCoreClient.make_async_mp_client(
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/tracing/otel.py", line 178, in sync_wrapper
    (APIServer pid=33938) return func(    args, kwargs)
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/core_client.py", line 130, in make_async_mp_client
    (APIServer pid=33938) return AsyncMPClient(client_args)
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/tracing/otel.py", line 178, in sync_wrapper
    (APIServer pid=33938) return func(    args,     kwargs)
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/core_client.py", line 900, in __init
    (APIServer pid=33938) super().init(
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/core_client.py", line 535, in init
    (APIServer pid=33938) with launch_core_engines(
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/contextlib.py", line 142, in exit
    (APIServer pid=33938) next(self.gen)
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/utils.py", line 1128, in launch_core_engines
    (APIServer pid=33938) wait_for_engine_startup(
    (APIServer pid=33938) File "/opt/conda/lib/python3.10/site-packages/vllm/v1/engine/utils.py", line 1187, in wait_for_engine_startup
    (APIServer pid=33938) raise RuntimeError(
    (APIServer pid=33938) RuntimeError: Engine core initialization failed. See root cause above. Failed core proc(s): {}
    root@18c870cbacae:/holfer/python_scripts# /opt/conda/lib/python3.10/multiprocessing/resource_tracker.py:224: UserWarning: resource_tracker: There appear to be 8 leaked shared_memory objects to clean up at shutdown
    warnings.warn('resource_tracker: There appear to be %d '

  • arrow_forward

    Thread has been moved from 公共.

  • link
    Members 794 posts
    2026年7月13日 10:05
    link

    尊敬的开发者您好，此镜像不支持DS V4 Flash

  • link
    Members 8 posts
    2026年7月13日 10:16
    link

    你使用的也是vllm0.21版本么？我部署你这个量化的flash也出现跟你一样的错误

  • arrow_forward

    Thread has been moved from 解决中.

  • link
    Members 2 posts
    2026年7月29日 14:12
    link

    那可以告诉我们，哪个镜像可以部署吗？

  • link
    Members 794 posts
    2026年7月29日 15:25
    link

    尊敬的开发者您好，已通过个人主题回复