root@10-66-1-235:/workspace# mkdir -p /data && python -m sglang.launch_server   --model-path /model/Qwen3-14B   --host 0.0.0.0   --enable-hierarchical-cache   --mem-fraction-static 0.55   --hicache-ratio 1   --hicache-storage-backend file   --hicache-storage-backend-extra-config '{"cache_dir": "/data"}'   --hicache-io-backend direct   --hicache-mem-layout page_first_direct   --page-size 64   --chunked-prefill-size 4096
INFO Print the version information of mcoplib during compilation.

Version info:Mcoplib_Version = '0.4.4'
Build_Maca_Version = '3.7.1.5'
GIT_BRANCH = 'HEAD'
GIT_COMMIT = '4d9c16e'
Vllm Op Version = 0.19.0
SGlang Op Version  = 0.5.10

INFO Staring Check the current MACA version of the operating environment.

INFO: Release major.minor matching,  successful:3.7.

W0706 09:59:25.968000 32 site-packages/torch/utils/cpp_extension.py:2527] TORCH_CUDA_ARCH_LIST is not set, all archs for visible cards are included for compilation.
W0706 09:59:25.968000 32 site-packages/torch/utils/cpp_extension.py:2527] If this is not desired, please set os.environ['TORCH_CUDA_ARCH_LIST'] to specific architectures.
/opt/conda/lib/python3.10/site-packages/torchvision/datapoints/__init__.py:12: UserWarning: The torchvision.datapoints and torchvision.transforms.v2 namespaces are still Beta. While we do not expect major breaking changes, some APIs may still change according to user feedback. Please submit any feedback you may have in this issue: https://github.com/pytorch/vision/issues/6753, and you can also check out https://github.com/pytorch/vision/issues/7319 to learn more about the APIs that we suspect might involve future changes. You can silence this warning by calling torchvision.disable_beta_transforms_warning().
  warnings.warn(_BETA_TRANSFORMS_WARNING)
/opt/conda/lib/python3.10/site-packages/torchvision/transforms/v2/__init__.py:54: UserWarning: The torchvision.datapoints and torchvision.transforms.v2 namespaces are still Beta. While we do not expect major breaking changes, some APIs may still change according to user feedback. Please submit any feedback you may have in this issue: https://github.com/pytorch/vision/issues/6753, and you can also check out https://github.com/pytorch/vision/issues/7319 to learn more about the APIs that we suspect might involve future changes. You can silence this warning by calling torchvision.disable_beta_transforms_warning().
  warnings.warn(_BETA_TRANSFORMS_WARNING)
/opt/conda/lib/python3.10/site-packages/sglang/launch_server.py:60: UserWarning: 'python -m sglang.launch_server' is still supported, but 'sglang serve' is the recommended entrypoint.
  Example: sglang serve --model-path <model> [options]
  warnings.warn(
[2026-07-06 09:59:29] Attention backend not specified. Use flashinfer backend by default.
[2026-07-06 09:59:29] Cuda graph max bs is adjusted to 128.
[2026-07-06 09:59:30] flashinfer.comm is not available, falling back to standard implementation
[2026-07-06 09:59:31] server_args=ServerArgs(model_path='/model/Qwen3-14B', tokenizer_path='/model/Qwen3-14B', tokenizer_mode='auto', tokenizer_backend='huggingface', tokenizer_worker_num=1, detokenizer_worker_num=1, skip_tokenizer_init=False, load_format='auto', model_loader_extra_config='{}', trust_remote_code=False, context_length=None, is_embedding=False, enable_multimodal=None, revision=None, model_impl='auto', host='0.0.0.0', port=30000, fastapi_root_path='', grpc_mode=False, skip_server_warmup=False, warmups=None, nccl_port=None, checkpoint_engine_wait_weights_before_ready=False, ssl_keyfile=None, ssl_certfile=None, ssl_ca_certs=None, ssl_keyfile_password=None, enable_ssl_refresh=False, enable_http2=False, dtype='auto', quantization=None, quantization_param_path=None, kv_cache_dtype='auto', enable_fp32_lm_head=False, modelopt_quant=None, modelopt_checkpoint_restore_path=None, modelopt_checkpoint_save_path=None, modelopt_export_path=None, quantize_and_serve=False, rl_quant_profile=None, mem_fraction_static=0.55, max_running_requests=None, max_queued_requests=None, max_total_tokens=None, chunked_prefill_size=4096, enable_dynamic_chunking=False, max_prefill_tokens=16384, prefill_max_requests=None, schedule_policy='fcfs', enable_priority_scheduling=False, disable_priority_preemption=False, default_priority_value=None, abort_on_priority_when_disabled=False, schedule_low_priority_values_first=False, priority_scheduling_preemption_threshold=10, schedule_conservativeness=1.0, page_size=64, swa_full_tokens_ratio=0.8, disable_hybrid_swa_memory=False, radix_eviction_policy='lru', enable_prefill_delayer=False, prefill_delayer_max_delay_passes=30, prefill_delayer_token_usage_low_watermark=None, prefill_delayer_forward_passes_buckets=None, prefill_delayer_wait_seconds_buckets=None, device='cuda', tp_size=1, pp_size=1, pp_max_micro_batch_size=None, pp_async_batch_depth=0, stream_interval=1, batch_notify_size=16, stream_response_default_include_usage=False, incremental_streaming_output=False, enable_streaming_session=False, random_seed=856938111, constrained_json_whitespace_pattern=None, constrained_json_disable_any_whitespace=False, watchdog_timeout=300, soft_watchdog_timeout=None, dist_timeout=None, download_dir=None, model_checksum=None, base_gpu_id=0, gpu_id_step=1, sleep_on_idle=False, use_ray=False, custom_sigquit_handler=None, log_level='info', log_level_http=None, log_requests=False, log_requests_level=2, log_requests_format='text', log_requests_target=None, uvicorn_access_log_exclude_prefixes=[], crash_dump_folder=None, show_time_cost=False, enable_metrics=False, grpc_http_sidecar_port=None, enable_mfu_metrics=False, enable_metrics_for_all_schedulers=False, tokenizer_metrics_custom_labels_header='x-custom-labels', tokenizer_metrics_allowed_custom_labels=None, extra_metric_labels=None, bucket_time_to_first_token=None, bucket_inter_token_latency=None, bucket_e2e_request_latency=None, prompt_tokens_buckets=None, generation_tokens_buckets=None, gc_warning_threshold_secs=0.0, decode_log_interval=40, enable_request_time_stats_logging=False, kv_events_config=None, enable_trace=False, otlp_traces_endpoint='localhost:4317', export_metrics_to_file=False, export_metrics_to_file_dir=None, api_key=None, admin_api_key=None, served_model_name='/model/Qwen3-14B', weight_version='default', chat_template=None, hf_chat_template_name=None, completion_template=None, file_storage_path='sglang_storage', enable_cache_report=False, reasoning_parser=None, strip_thinking_cache=False, tool_call_parser=None, tool_server=None, sampling_defaults='model', dp_size=1, load_balance_method='round_robin', attn_cp_size=1, moe_dp_size=1, dist_init_addr=None, nnodes=1, node_rank=0, json_model_override_args='{}', preferred_sampling_params=None, enable_lora=None, enable_lora_overlap_loading=None, max_lora_rank=None, lora_target_modules=None, lora_paths=None, max_loaded_loras=None, max_loras_per_batch=8, lora_eviction_policy='lru', lora_backend='csgmv', max_lora_chunk_size=16, experts_shared_outer_loras=None, lora_use_virtual_experts=False, lora_strict_loading=False, attention_backend='flashinfer', decode_attention_backend=None, prefill_attention_backend=None, sampling_backend='flashinfer', grammar_backend='xgrammar', mm_attention_backend=None, fp8_gemm_runner_backend='auto', fp4_gemm_runner_backend='auto', nsa_prefill_backend='flashmla_sparse', nsa_decode_backend='flashmla_kv', disable_flashinfer_autotune=False, mamba_backend='triton', speculative_algorithm=None, speculative_draft_model_path=None, speculative_draft_model_revision=None, speculative_draft_load_format=None, speculative_num_steps=None, speculative_eagle_topk=None, speculative_num_draft_tokens=None, speculative_dflash_block_size=None, speculative_dflash_draft_window_size=None, speculative_accept_threshold_single=1.0, speculative_accept_threshold_acc=1.0, speculative_token_map=None, speculative_attention_mode='prefill', speculative_draft_attention_backend=None, speculative_moe_runner_backend='auto', speculative_moe_a2a_backend=None, speculative_draft_model_quantization=None, speculative_adaptive=False, speculative_adaptive_config=None, speculative_skip_dp_mlp_sync=False, speculative_ngram_min_bfs_breadth=1, speculative_ngram_max_bfs_breadth=10, speculative_ngram_match_type='BFS', speculative_ngram_max_trie_depth=18, speculative_ngram_capacity=10000000, speculative_ngram_external_corpus_path=None, speculative_ngram_external_sam_budget=0, speculative_ngram_external_corpus_max_tokens=10000000, enable_multi_layer_eagle=False, ep_size=1, moe_a2a_backend='none', moe_runner_backend='auto', record_nolora_graph=True, flashinfer_mxfp4_moe_precision='default', enable_flashinfer_allreduce_fusion=False, enforce_disable_flashinfer_allreduce_fusion=False, enable_aiter_allreduce_fusion=False, deepep_mode='auto', ep_num_redundant_experts=0, ep_dispatch_algorithm=None, init_expert_location='trivial', enable_eplb=False, eplb_algorithm='auto', eplb_rebalance_num_iterations=1000, eplb_rebalance_layers_per_chunk=None, eplb_min_rebalancing_utilization_threshold=1.0, expert_distribution_recorder_mode=None, expert_distribution_recorder_buffer_size=1000, enable_expert_distribution_metrics=False, deepep_config=None, moe_dense_tp_size=None, elastic_ep_backend=None, enable_elastic_expert_backup=False, mooncake_ib_device=None, elastic_ep_rejoin=False, max_mamba_cache_size=None, mamba_ssm_dtype=None, mamba_full_memory_ratio=0.9, mamba_scheduler_strategy='no_buffer', mamba_track_interval=256, linear_attn_backend='triton', linear_attn_decode_backend=None, linear_attn_prefill_backend=None, enable_hierarchical_cache=True, hicache_ratio=1.0, hicache_size=0, hicache_write_policy='write_through', hicache_io_backend='direct', hicache_mem_layout='page_first_direct', hicache_storage_backend='file', hicache_storage_prefetch_policy='best_effort', hicache_storage_backend_extra_config='{"cache_dir": "/data"}', enable_hisparse=False, hisparse_config=None, enable_lmcache=False, kt_weight_path=None, kt_method='AMXINT4', kt_cpuinfer=None, kt_threadpool_count=2, kt_num_gpu_experts=None, kt_max_deferred_experts_per_token=None, dllm_algorithm=None, dllm_algorithm_config=None, cpu_offload_gb=0, offload_group_size=-1, offload_num_in_group=1, offload_prefetch_step=1, offload_mode='cpu', enable_mis=False, disable_radix_cache=False, cuda_graph_max_bs=128, cuda_graph_bs=None, disable_cuda_graph=False, disable_cuda_graph_padding=False, enable_breakable_cuda_graph=False, enable_profile_cuda_graph=False, enable_cudagraph_gc=False, debug_cuda_graph=False, enable_layerwise_nvtx_marker=False, enable_nccl_nvls=False, enable_symm_mem=False, disable_flashinfer_cutlass_moe_fp4_allgather=False, enable_tokenizer_batch_encode=False, disable_tokenizer_batch_decode=False, disable_outlines_disk_cache=False, disable_custom_all_reduce=False, enable_mscclpp=False, enable_torch_symm_mem=False, pre_warm_nccl=False, disable_overlap_schedule=False, enable_mixed_chunk=False, enable_dp_attention=False, enable_dp_attention_local_control_broadcast=False, enable_dp_lm_head=False, enable_two_batch_overlap=False, enable_single_batch_overlap=False, tbo_token_distribution_threshold=0.48, enable_torch_compile=False, disable_piecewise_cuda_graph=True, enforce_piecewise_cuda_graph=False, enable_torch_compile_debug_mode=False, torch_compile_max_bs=32, piecewise_cuda_graph_max_tokens=4096, piecewise_cuda_graph_tokens=[4, 8, 12, 16, 20, 24, 28, 32, 48, 64, 80, 96, 112, 128, 144, 160, 176, 192, 208, 224, 240, 256, 288, 320, 352, 384, 416, 448, 480, 512, 576, 640, 704, 768, 832, 896, 960, 1024, 1280, 1536, 1792, 2048, 2304, 2560, 2816, 3072, 3328, 3584, 3840, 4096], piecewise_cuda_graph_compiler='eager', torchao_config='', enable_nan_detection=False, enable_p2p_check=False, triton_attention_reduce_in_fp32=False, triton_attention_num_kv_splits=8, triton_attention_split_tile_size=None, num_continuous_decode_steps=1, delete_ckpt_after_loading=False, enable_memory_saver=False, enable_weights_cpu_backup=False, enable_draft_weights_cpu_backup=False, allow_auto_truncate=False, enable_custom_logit_processor=False, flashinfer_mla_disable_ragged=False, disable_shared_experts_fusion=False, enforce_shared_experts_fusion=False, disable_chunked_prefix_cache=False, disable_fast_image_processor=False, keep_mm_feature_on_device=False, enable_return_hidden_states=False, enable_return_routed_experts=False, scheduler_recv_interval=1, numa_node=None, enable_deterministic_inference=False, rl_on_policy_target=None, enable_attn_tp_input_scattered=False, gc_threshold=None, enable_nsa_prefill_context_parallel=False, nsa_prefill_cp_mode='round-robin-split', enable_fused_qk_norm_rope=False, enable_precise_embedding_interpolation=False, enable_fused_moe_sum_all_reduce=False, enable_prefill_context_parallel=False, prefill_cp_mode='in-seq-split', enable_dynamic_batch_tokenizer=False, dynamic_batch_tokenizer_batch_size=32, dynamic_batch_tokenizer_batch_timeout=0.002, debug_tensor_dump_output_folder=None, debug_tensor_dump_layers=None, debug_tensor_dump_input_file=None, debug_tensor_dump_inject=False, disaggregation_mode='null', disaggregation_transfer_backend='mooncake', disaggregation_bootstrap_port=8998, disaggregation_ib_device=None, disaggregation_decode_enable_radix_cache=False, disaggregation_decode_enable_offload_kvcache=False, num_reserved_decode_tokens=512, disaggregation_decode_polling_interval=1, encoder_only=False, language_only=False, encoder_transfer_backend='zmq_to_scheduler', encoder_urls=[], enable_adaptive_dispatch_to_encoder=False, custom_weight_loader=[], weight_loader_disable_mmap=False, weight_loader_prefetch_checkpoints=False, weight_loader_prefetch_num_threads=4, remote_instance_weight_loader_seed_instance_ip=None, remote_instance_weight_loader_seed_instance_service_port=None, remote_instance_weight_loader_send_weights_group_ports=None, remote_instance_weight_loader_backend='nccl', remote_instance_weight_loader_start_seed_via_transfer_engine=False, engine_info_bootstrap_port=6789, modelexpress_config=None, enable_pdmux=False, pdmux_config_path=None, sm_group_num=8, enable_broadcast_mm_inputs_process=False, enable_prefix_mm_cache=False, mm_enable_dp_encoder=False, mm_process_config={}, limit_mm_data_per_request=None, enable_mm_global_cache=False, decrypted_config_file=None, decrypted_draft_config_file=None, forward_hooks=None, enable_quant_communications=False, msprobe_dump_config=None)
[2026-07-06 09:59:32] Using default HuggingFace chat template with detected content format: string
INFO Print the version information of mcoplib during compilation.

Version info:Mcoplib_Version = '0.4.4'
Build_Maca_Version = '3.7.1.5'
GIT_BRANCH = 'HEAD'
GIT_COMMIT = '4d9c16e'
Vllm Op Version = 0.19.0
SGlang Op Version  = 0.5.10

INFO Staring Check the current MACA version of the operating environment.

INFO: Release major.minor matching,  successful:3.7.

INFO Print the version information of mcoplib during compilation.

Version info:Mcoplib_Version = '0.4.4'
Build_Maca_Version = '3.7.1.5'
GIT_BRANCH = 'HEAD'
GIT_COMMIT = '4d9c16e'
Vllm Op Version = 0.19.0
SGlang Op Version  = 0.5.10

INFO Staring Check the current MACA version of the operating environment.

INFO: Release major.minor matching,  successful:3.7.

W0706 09:59:51.975000 244 site-packages/torch/utils/cpp_extension.py:2527] TORCH_CUDA_ARCH_LIST is not set, all archs for visible cards are included for compilation.
W0706 09:59:51.975000 244 site-packages/torch/utils/cpp_extension.py:2527] If this is not desired, please set os.environ['TORCH_CUDA_ARCH_LIST'] to specific architectures.
W0706 09:59:52.184000 243 site-packages/torch/utils/cpp_extension.py:2527] TORCH_CUDA_ARCH_LIST is not set, all archs for visible cards are included for compilation.
W0706 09:59:52.184000 243 site-packages/torch/utils/cpp_extension.py:2527] If this is not desired, please set os.environ['TORCH_CUDA_ARCH_LIST'] to specific architectures.
/opt/conda/lib/python3.10/site-packages/torchvision/datapoints/__init__.py:12: UserWarning: The torchvision.datapoints and torchvision.transforms.v2 namespaces are still Beta. While we do not expect major breaking changes, some APIs may still change according to user feedback. Please submit any feedback you may have in this issue: https://github.com/pytorch/vision/issues/6753, and you can also check out https://github.com/pytorch/vision/issues/7319 to learn more about the APIs that we suspect might involve future changes. You can silence this warning by calling torchvision.disable_beta_transforms_warning().
  warnings.warn(_BETA_TRANSFORMS_WARNING)
/opt/conda/lib/python3.10/site-packages/torchvision/transforms/v2/__init__.py:54: UserWarning: The torchvision.datapoints and torchvision.transforms.v2 namespaces are still Beta. While we do not expect major breaking changes, some APIs may still change according to user feedback. Please submit any feedback you may have in this issue: https://github.com/pytorch/vision/issues/6753, and you can also check out https://github.com/pytorch/vision/issues/7319 to learn more about the APIs that we suspect might involve future changes. You can silence this warning by calling torchvision.disable_beta_transforms_warning().
  warnings.warn(_BETA_TRANSFORMS_WARNING)
/opt/conda/lib/python3.10/site-packages/torchvision/datapoints/__init__.py:12: UserWarning: The torchvision.datapoints and torchvision.transforms.v2 namespaces are still Beta. While we do not expect major breaking changes, some APIs may still change according to user feedback. Please submit any feedback you may have in this issue: https://github.com/pytorch/vision/issues/6753, and you can also check out https://github.com/pytorch/vision/issues/7319 to learn more about the APIs that we suspect might involve future changes. You can silence this warning by calling torchvision.disable_beta_transforms_warning().
  warnings.warn(_BETA_TRANSFORMS_WARNING)
/opt/conda/lib/python3.10/site-packages/torchvision/transforms/v2/__init__.py:54: UserWarning: The torchvision.datapoints and torchvision.transforms.v2 namespaces are still Beta. While we do not expect major breaking changes, some APIs may still change according to user feedback. Please submit any feedback you may have in this issue: https://github.com/pytorch/vision/issues/6753, and you can also check out https://github.com/pytorch/vision/issues/7319 to learn more about the APIs that we suspect might involve future changes. You can silence this warning by calling torchvision.disable_beta_transforms_warning().
  warnings.warn(_BETA_TRANSFORMS_WARNING)
flashinfer.comm is not available, falling back to standard implementation
flashinfer.comm is not available, falling back to standard implementation
[2026-07-06 09:59:57] Init torch distributed begin.
[Gloo] Rank 0 is connected to 0 peer ranks. Expected number of connected peer ranks is : 0
[Gloo] Rank 0 is connected to 0 peer ranks. Expected number of connected peer ranks is : 0
[2026-07-06 09:59:57] DCP disabled, dcp_size=1, tp_size=1
[Gloo] Rank 0 is connected to 0 peer ranks. Expected number of connected peer ranks is : 0
[Gloo] Rank 0 is connected to 0 peer ranks. Expected number of connected peer ranks is : 0
[2026-07-06 09:59:57] Init torch distributed ends. elapsed=0.06 s, mem usage=0.06 GB
[2026-07-06 09:59:58] Load weight begin. avail mem=61.91 GB
Multi-thread loading shards: 100% Completed | 8/8 [00:13<00:00,  1.70s/it]
[2026-07-06 10:00:12] Load weight end. elapsed=13.73 s, type=Qwen3ForCausalLM, avail mem=28.91 GB, mem usage=33.00 GB.
[2026-07-06 10:00:12] Using KV cache dtype: torch.bfloat16
[2026-07-06 10:00:12] KV Cache is allocated. #tokens: 6848, K size: 0.53 GB, V size: 0.53 GB
[2026-07-06 10:00:12] Memory pool end. avail mem=27.91 GB
[2026-07-06 10:00:13] Capture cuda graph begin. This can take up to several minutes. avail mem=27.10 GB
[2026-07-06 10:00:13] Capture cuda graph bs [1, 2, 4, 8, 16, 24, 32, 40, 48, 56, 64, 72, 80, 88, 96, 104, 112, 120, 128]
Capturing batches (bs=128 avail_mem=26.97 GB):   0%|                                                                                                                                                             | 0/19 [00:00<?, ?it/s][rank0]:W0706 10:00:14.104000 243 site-packages/torch/utils/cpp_extension.py:2527] TORCH_CUDA_ARCH_LIST is not set, all archs for visible cards are included for compilation.
[rank0]:W0706 10:00:14.104000 243 site-packages/torch/utils/cpp_extension.py:2527] If this is not desired, please set os.environ['TORCH_CUDA_ARCH_LIST'] to specific architectures.
[rank0]:W0706 10:00:14.105000 243 site-packages/torch/utils/cpp_extension.py:2527] TORCH_CUDA_ARCH_LIST is not set, all archs for visible cards are included for compilation.
[rank0]:W0706 10:00:14.105000 243 site-packages/torch/utils/cpp_extension.py:2527] If this is not desired, please set os.environ['TORCH_CUDA_ARCH_LIST'] to specific architectures.
[rank0]:W0706 10:00:14.110000 243 site-packages/torch/utils/cpp_extension.py:2527] TORCH_CUDA_ARCH_LIST is not set, all archs for visible cards are included for compilation.
[rank0]:W0706 10:00:14.110000 243 site-packages/torch/utils/cpp_extension.py:2527] If this is not desired, please set os.environ['TORCH_CUDA_ARCH_LIST'] to specific architectures.
[rank0]:W0706 10:00:14.111000 243 site-packages/torch/utils/cpp_extension.py:2527] TORCH_CUDA_ARCH_LIST is not set, all archs for visible cards are included for compilation.
[rank0]:W0706 10:00:14.111000 243 site-packages/torch/utils/cpp_extension.py:2527] If this is not desired, please set os.environ['TORCH_CUDA_ARCH_LIST'] to specific architectures.
Capturing batches (bs=1 avail_mem=26.59 GB): 100%|██████████████████████████████████████████████████████████████████████████████████████████████████████████████████████████████████████████████████████| 19/19 [00:32<00:00,  1.70s/it]
[2026-07-06 10:00:45] Capture cuda graph end. Time elapsed: 32.92 s. mem usage=0.56 GB. avail mem=26.53 GB.
[2026-07-06 10:00:45] Disable piecewise CUDA graph because --disable-piecewise-cuda-graph is set
[2026-07-06 10:00:46] max_total_num_tokens=6848, chunked_prefill_size=4096, max_prefill_tokens=16384, max_running_requests=2048, context_len=40960, available_gpu_mem=26.53 GB
[2026-07-06 10:00:46] Allocating 1.13 GB host memory for hierarchical KV cache.
[2026-07-06 10:00:48] Failed to load JIT HiCache kernel: ninja exited with status 1
stdout:
[1/2] /opt/maca/tools/cu-bridge/bin/nvcc --generate-dependencies-with-compile --dependency-output cuda_0.o.d -Xcompiler -fPIC -std=c++17 -O2 -gencode=arch=compute_80,code=sm_80 -DSGL_CUDA_ARCH=800 -std=c++20 -O3 --expt-relaxed-constexpr -I/opt/conda/lib/python3.10/site-packages/tvm_ffi/include -I/opt/conda/lib/python3.10/site-packages/tvm_ffi/include -I/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/include -c /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu -o cuda_0.o
FAILED: cuda_0.o
/opt/maca/tools/cu-bridge/bin/nvcc --generate-dependencies-with-compile --dependency-output cuda_0.o.d -Xcompiler -fPIC -std=c++17 -O2 -gencode=arch=compute_80,code=sm_80 -DSGL_CUDA_ARCH=800 -std=c++20 -O3 --expt-relaxed-constexpr -I/opt/conda/lib/python3.10/site-packages/tvm_ffi/include -I/opt/conda/lib/python3.10/site-packages/tvm_ffi/include -I/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/include -c /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu -o cuda_0.o
In file included from /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:7:
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:40:71: error: invalid input constraint 'l' in asm
   40 |   asm volatile("ld.global.L1::no_allocate.b32 %0,[%1];" : "=r"(tmp) : "l"(src));
      |                                                                       ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:46:92: error: invalid input constraint 'l' in asm
   46 |   asm volatile("ld.global.L1::no_allocate.v2.b32 {%0,%1},[%2];" : "=r"(tmp0), "=r"(tmp1) : "l"(src));
      |                                                                                            ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:54:18: error: invalid input constraint 'l' in asm
   54 |                : "l"(src));
      |                  ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:60:59: error: invalid input constraint 'l' in asm
   60 |   asm volatile("st.global.L1::no_allocate.b32 [%0],%1;" ::"l"(dst), "r"(tmp));
      |                                                           ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:66:67: error: invalid input constraint 'l' in asm
   66 |   asm volatile("st.global.L1::no_allocate.v2.b32 [%0],{%1,%2};" ::"l"(dst), "r"(tmp0), "r"(tmp1));
      |                                                                   ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:75:64: error: invalid input constraint 'l' in asm
   75 |       "st.global.L1::no_allocate.v4.b32 [%0],{%1,%2,%3,%4};" ::"l"(dst), "r"(tmp0), "r"(tmp1), "r"(tmp2), "r"(tmp3));
      |                                                                ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:144:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  144 | SGL_HICACHE_KERNEL void hicache_transfer_per_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:183:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  183 | SGL_HICACHE_KERNEL void hicache_transfer_all_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:144:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  144 | SGL_HICACHE_KERNEL void hicache_transfer_per_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:227:38: note: in instantiation of function template specialization '(anonymous namespace)::hicache_transfer_per_layer<int, 2048L, 1U, 2U, 1024U, false>' requested here
  227 |   static constexpr auto kernel_one = hicache_transfer_per_layer<T, kElementSize, kUnroll, kBlockQuota, kBlockSize>;
      |                                      ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:302:37: note: in instantiation of static data member '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::kernel_one<int>' requested here
  302 |     const auto kernel = use_int32 ? kernel_one<int32_t> : kernel_one<int64_t>;
      |                                     ^
/root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:8:78: note: in instantiation of member function '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::run_one' requested here
    8 | TVM_FFI_DLL_EXPORT_TYPED_FUNC(launch_one, (&HiCacheKernel<2048, 1, 2, 1024>::run_one));
      |                                                                              ^
In file included from /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:7:
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:144:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  144 | SGL_HICACHE_KERNEL void hicache_transfer_per_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:227:38: note: in instantiation of function template specialization '(anonymous namespace)::hicache_transfer_per_layer<long, 2048L, 1U, 2U, 1024U, false>' requested here
  227 |   static constexpr auto kernel_one = hicache_transfer_per_layer<T, kElementSize, kUnroll, kBlockQuota, kBlockSize>;
      |                                      ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:302:59: note: in instantiation of static data member '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::kernel_one<long>' requested here
  302 |     const auto kernel = use_int32 ? kernel_one<int32_t> : kernel_one<int64_t>;
      |                                                           ^
/root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:8:78: note: in instantiation of member function '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::run_one' requested here
    8 | TVM_FFI_DLL_EXPORT_TYPED_FUNC(launch_one, (&HiCacheKernel<2048, 1, 2, 1024>::run_one));
      |                                                                              ^
In file included from /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:7:
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:183:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  183 | SGL_HICACHE_KERNEL void hicache_transfer_all_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:229:38: note: in instantiation of function template specialization '(anonymous namespace)::hicache_transfer_all_layer<int, 2048L, 1U, 2U, 1024U, false>' requested here
  229 |   static constexpr auto kernel_all = hicache_transfer_all_layer<T, kElementSize, kUnroll, kBlockQuota, kBlockSize>;
      |                                      ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:360:37: note: in instantiation of static data member '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::kernel_all<int>' requested here
  360 |     const auto kernel = use_int32 ? kernel_all<int32_t> : kernel_all<int64_t>;
      |                                     ^
/root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:9:78: note: in instantiation of member function '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::run_all' requested here
    9 | TVM_FFI_DLL_EXPORT_TYPED_FUNC(launch_all, (&HiCacheKernel<2048, 1, 2, 1024>::run_all));
      |                                                                              ^
In file included from /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:7:
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:183:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  183 | SGL_HICACHE_KERNEL void hicache_transfer_all_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:229:38: note: in instantiation of function template specialization '(anonymous namespace)::hicache_transfer_all_layer<long, 2048L, 1U, 2U, 1024U, false>' requested here
  229 |   static constexpr auto kernel_all = hicache_transfer_all_layer<T, kElementSize, kUnroll, kBlockQuota, kBlockSize>;
      |                                      ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:360:59: note: in instantiation of static data member '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::kernel_all<long>' requested here
  360 |     const auto kernel = use_int32 ? kernel_all<int32_t> : kernel_all<int64_t>;
      |                                                           ^
/root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:9:78: note: in instantiation of member function '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::run_all' requested here
    9 | TVM_FFI_DLL_EXPORT_TYPED_FUNC(launch_all, (&HiCacheKernel<2048, 1, 2, 1024>::run_all));
      |                                                                              ^
In file included from /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:7:
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:144:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  144 | SGL_HICACHE_KERNEL void hicache_transfer_per_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:232:7: note: in instantiation of function template specialization '(anonymous namespace)::hicache_transfer_per_layer<int, 2048L, 1U, 2U, 1024U, true>' requested here
  232 |       hicache_transfer_per_layer<T, kElementSize, kUnroll, kBlockQuota, kBlockSize, true>;
      |       ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:422:37: note: in instantiation of static data member '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::kernel_one_mla<int>' requested here
  422 |     const auto kernel = use_int32 ? kernel_one_mla<int32_t> : kernel_one_mla<int64_t>;
      |                                     ^
/root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:10:82: note: in instantiation of member function '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::run_one_mla' requested here
   10 | TVM_FFI_DLL_EXPORT_TYPED_FUNC(launch_one_mla, (&HiCacheKernel<2048, 1, 2, 1024>::run_one_mla));
      |                                                                                  ^
In file included from /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:7:
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:144:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  144 | SGL_HICACHE_KERNEL void hicache_transfer_per_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:232:7: note: in instantiation of function template specialization '(anonymous namespace)::hicache_transfer_per_layer<long, 2048L, 1U, 2U, 1024U, true>' requested here
  232 |       hicache_transfer_per_layer<T, kElementSize, kUnroll, kBlockQuota, kBlockSize, true>;
      |       ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:422:63: note: in instantiation of static data member '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::kernel_one_mla<long>' requested here
  422 |     const auto kernel = use_int32 ? kernel_one_mla<int32_t> : kernel_one_mla<int64_t>;
      |                                                               ^
/root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:10:82: note: in instantiation of member function '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::run_one_mla' requested here
   10 | TVM_FFI_DLL_EXPORT_TYPED_FUNC(launch_one_mla, (&HiCacheKernel<2048, 1, 2, 1024>::run_one_mla));
      |                                                                                  ^
In file included from /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:7:
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:183:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  183 | SGL_HICACHE_KERNEL void hicache_transfer_all_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:235:7: note: in instantiation of function template specialization '(anonymous namespace)::hicache_transfer_all_layer<int, 2048L, 1U, 2U, 1024U, true>' requested here
  235 |       hicache_transfer_all_layer<T, kElementSize, kUnroll, kBlockQuota, kBlockSize, true>;
      |       ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:473:37: note: in instantiation of static data member '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::kernel_all_mla<int>' requested here
  473 |     const auto kernel = use_int32 ? kernel_all_mla<int32_t> : kernel_all_mla<int64_t>;
      |                                     ^
/root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:11:82: note: in instantiation of member function '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::run_all_mla' requested here
   11 | TVM_FFI_DLL_EXPORT_TYPED_FUNC(launch_all_mla, (&HiCacheKernel<2048, 1, 2, 1024>::run_all_mla));
      |                                                                                  ^
In file included from /root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:7:
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:183:1: warning: set minimum blocks' number is illegal in Maca and the value will be ignored [-Wmaca-min-blocks-per-multiprocessor]
  183 | SGL_HICACHE_KERNEL void hicache_transfer_all_layer(const  HicacheKernelParams params) {
      | ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:122:69: note: expanded from macro 'SGL_HICACHE_KERNEL'
  122 | #define SGL_HICACHE_KERNEL __global__ __launch_bounds__(kBlockSize, 1)
      |                                                                     ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:235:7: note: in instantiation of function template specialization '(anonymous namespace)::hicache_transfer_all_layer<long, 2048L, 1U, 2U, 1024U, true>' requested here
  235 |       hicache_transfer_all_layer<T, kElementSize, kUnroll, kBlockQuota, kBlockSize, true>;
      |       ^
/opt/conda/lib/python3.10/site-packages/sglang/jit_kernel/csrc/hicache.cuh:473:63: note: in instantiation of static data member '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::kernel_all_mla<long>' requested here
  473 |     const auto kernel = use_int32 ? kernel_all_mla<int32_t> : kernel_all_mla<int64_t>;
      |                                                               ^
/root/.cache/tvm-ffi/sgl_kernel_jit_hicache_2048_1_2_1024_e14cbc20997ec50a/cuda.cu:11:82: note: in instantiation of member function '(anonymous namespace)::HiCacheKernel<2048, 1, 2, 1024>::run_all_mla' requested here
   11 | TVM_FFI_DLL_EXPORT_TYPED_FUNC(launch_all_mla, (&HiCacheKernel<2048, 1, 2, 1024>::run_all_mla));
      |                                                                                  ^
10 warnings and 6 errors generated when compiling for xcore1000.
ninja: build stopped: subcommand failed.

[2026-07-06 10:00:48] Creating storage backend 'file' (sglang.srt.mem_cache.hicache_storage.HiCacheFile)
[2026-07-06 10:00:48] Created HiCacheFile storage directory at /tmp/hicache
[2026-07-06 10:00:50] INFO:     Started server process [32]
[2026-07-06 10:00:50] INFO:     Waiting for application startup.
[2026-07-06 10:00:50] Using default chat sampling params from model generation config: {'temperature': 0.6, 'top_k': 20, 'top_p': 0.95}
[2026-07-06 10:00:50] INFO:     Application startup complete.
[2026-07-06 10:00:50] INFO:     Uvicorn running on http://0.0.0.0:30000 (Press CTRL+C to quit)
[2026-07-06 10:00:51] INFO:     127.0.0.1:44530 - "GET /model_info HTTP/1.1" 200 OK
[2026-07-06 10:01:05] Prefill batch, #new-seq: 1, #new-token: 64, #cached-token: 0, token usage: 0.01, #running-req: 0, #queue-req: 0, #pending-token: 0, cuda graph: False, input throughput (token/s): 0.93
[2026-07-06 10:01:05] INFO:     127.0.0.1:44532 - "POST /generate HTTP/1.1" 200 OK
[2026-07-06 10:01:05] The server is fired up and ready to roll!