chore: cleanup garbage code (#29770)
This commit is contained in:
@@ -52,7 +52,6 @@ class CustomOp(nn.Module):
|
||||
def forward_tpu(self, *args, **kwargs) -> Any:
|
||||
# By default, we assume that TPU ops are compatible with the
|
||||
# PyTorch-native implementation.
|
||||
# NOTE(woosuk): This is a placeholder for future extensions.
|
||||
return self.forward_native(*args, **kwargs)
|
||||
|
||||
def forward_musa(self, *args, **kwargs) -> Any:
|
||||
|
||||
@@ -854,7 +854,7 @@ class OpenAIServingChat(OpenAIServingBase):
|
||||
prompt_ids = self.tokenizer_manager.tokenizer.encode(
|
||||
rendered_prompt, **encode_kwargs
|
||||
)
|
||||
except Exception as e:
|
||||
except Exception:
|
||||
# If the first attempt fails, try with flat function-only format.
|
||||
# Some templates (e.g. Mistral) expect tools without the OpenAI wrapper.
|
||||
tools = (
|
||||
|
||||
@@ -1689,7 +1689,7 @@ class OpenAIServingResponses(OpenAIServingChat):
|
||||
)
|
||||
|
||||
async def empty_async_generator():
|
||||
if False:
|
||||
for _ in ():
|
||||
yield
|
||||
|
||||
final_response = await self.responses_full_generator(
|
||||
|
||||
@@ -801,13 +801,9 @@ class _DetailAccumulator(_UtilizationRateAccumulatorMixin):
|
||||
self._records = []
|
||||
|
||||
def get_single_pass_gatherer_keys(self):
|
||||
if False: # TODO `server_args.enable_two_batch_overlap`
|
||||
return [_SINGLE_PASS_GATHERER_KEY_PRIMARY, "child_a", "child_b"]
|
||||
return super().get_single_pass_gatherer_keys()
|
||||
|
||||
def get_single_pass_gatherer_key(self, debug_name: Optional[str]):
|
||||
if False: # TODO `server_args.enable_two_batch_overlap`
|
||||
return debug_name or _SINGLE_PASS_GATHERER_KEY_PRIMARY
|
||||
return super().get_single_pass_gatherer_key(debug_name)
|
||||
|
||||
def append(
|
||||
|
||||
@@ -280,14 +280,8 @@ class XIELU(MultiPlatformOp):
|
||||
)
|
||||
self._xielu_cuda_fn = self._xielu_cuda
|
||||
logger.warning_once(msg)
|
||||
except Exception as err:
|
||||
except Exception:
|
||||
pass
|
||||
# logger.warning_once(
|
||||
# "CUDA-fused xIELU not available (%s) –"
|
||||
# " falling back to a Python version.\n"
|
||||
# "For CUDA xIELU (experimental), `pip install git+https://github.com/nickjbrowning/XIELU`",
|
||||
# str(err),
|
||||
# )
|
||||
|
||||
def _xielu_python(self, x: torch.Tensor) -> torch.Tensor:
|
||||
alpha_p = nn.functional.softplus(self.alpha_p)
|
||||
|
||||
@@ -253,135 +253,11 @@ class GetKAndS:
|
||||
)
|
||||
|
||||
|
||||
class SetK:
|
||||
@classmethod
|
||||
def execute(cls, *args, buf, **kwargs):
|
||||
return cls.torch_fast(*args, **kwargs, buf=buf)
|
||||
|
||||
@classmethod
|
||||
def slow(
|
||||
cls,
|
||||
pool: "DSATokenToKVPool",
|
||||
buf: torch.Tensor,
|
||||
loc: torch.Tensor,
|
||||
index_k: torch.Tensor,
|
||||
):
|
||||
for i in range(len(loc)):
|
||||
page_index = loc[i] // pool.page_size
|
||||
offset = loc[i] % pool.page_size
|
||||
buf[
|
||||
page_index,
|
||||
offset * pool.index_head_dim : (offset + 1) * pool.index_head_dim,
|
||||
] = index_k[i].view(torch.uint8)
|
||||
|
||||
@classmethod
|
||||
def torch_fast(
|
||||
cls,
|
||||
pool: "DSATokenToKVPool",
|
||||
buf: torch.Tensor,
|
||||
loc: torch.Tensor,
|
||||
index_k: torch.Tensor,
|
||||
):
|
||||
(num_tokens_to_write,) = loc.shape
|
||||
buf_numel_per_page = buf.shape[1]
|
||||
num_k_bytes_per_token = pool.index_head_dim
|
||||
|
||||
# loc: (num_tokens_to_write,), int32, element := the token index to write to
|
||||
loc_page_index = loc // pool.page_size
|
||||
loc_token_offset_in_page = loc % pool.page_size
|
||||
|
||||
flat_buf = buf.flatten()
|
||||
flat_indices = (
|
||||
(loc_page_index * buf_numel_per_page)[:, None]
|
||||
+ (loc_token_offset_in_page * num_k_bytes_per_token)[:, None]
|
||||
+ torch.arange(num_k_bytes_per_token, dtype=torch.int32, device="cuda")[
|
||||
None, :
|
||||
]
|
||||
)
|
||||
num_k_bytes_total = num_tokens_to_write * num_k_bytes_per_token
|
||||
flat_indices = flat_indices.flatten()[:num_k_bytes_total]
|
||||
flat_buf[flat_indices] = index_k.view(torch.uint8).flatten()
|
||||
|
||||
|
||||
class SetS:
|
||||
@classmethod
|
||||
def execute(cls, *args, buf, **kwargs):
|
||||
return cls.torch_fast(*args, **kwargs, buf=buf)
|
||||
|
||||
@classmethod
|
||||
def slow(
|
||||
cls,
|
||||
pool: "DSATokenToKVPool",
|
||||
buf: torch.Tensor,
|
||||
loc: torch.Tensor,
|
||||
index_k_scale: torch.Tensor,
|
||||
):
|
||||
for i in range(len(loc)):
|
||||
page_index = loc[i] // pool.page_size
|
||||
offset = loc[i] % pool.page_size
|
||||
start = pool.page_size * pool.index_head_dim
|
||||
buf[page_index, start + offset * 4 : start + (offset + 1) * 4] = (
|
||||
index_k_scale[i].view(torch.uint8)
|
||||
)
|
||||
|
||||
@classmethod
|
||||
def torch_fast(
|
||||
cls,
|
||||
pool: "DSATokenToKVPool",
|
||||
buf: torch.Tensor,
|
||||
loc: torch.Tensor,
|
||||
index_k_scale: torch.Tensor,
|
||||
):
|
||||
(num_tokens_to_write,) = loc.shape
|
||||
buf_numel_per_page = buf.shape[1]
|
||||
num_s_bytes_per_token = 4
|
||||
s_offset_in_page = pool.page_size * pool.index_head_dim
|
||||
|
||||
# loc: (num_tokens_to_write,), int32, element := the token index to write to
|
||||
loc_page_index = loc // pool.page_size
|
||||
loc_token_offset_in_page = loc % pool.page_size
|
||||
|
||||
flat_buf = buf.flatten()
|
||||
flat_indices = (
|
||||
(loc_page_index * buf_numel_per_page)[:, None]
|
||||
+ s_offset_in_page
|
||||
+ (loc_token_offset_in_page * num_s_bytes_per_token)[:, None]
|
||||
+ torch.arange(num_s_bytes_per_token, dtype=torch.int32, device="cuda")[
|
||||
None, :
|
||||
]
|
||||
)
|
||||
number_s_bytes_total = num_tokens_to_write * num_s_bytes_per_token
|
||||
flat_indices = flat_indices.flatten()[:number_s_bytes_total]
|
||||
flat_buf[flat_indices] = index_k_scale.view(torch.uint8).flatten()
|
||||
|
||||
|
||||
class SetKAndS:
|
||||
@classmethod
|
||||
def execute(cls, *args, buf, **kwargs):
|
||||
if 0:
|
||||
# print("SetK, SetS comparison test")
|
||||
buf_cloned = buf.clone()
|
||||
cls.vanilla(*args, **kwargs, buf=buf)
|
||||
cls.triton(*args, **kwargs, buf=buf_cloned)
|
||||
|
||||
def _clear_token_0(target):
|
||||
target[0, :128] = target[0, 64 * 128 : 64 * 128 + 4] = 0
|
||||
|
||||
_clear_token_0(buf)
|
||||
_clear_token_0(buf_cloned)
|
||||
|
||||
assert torch.all(
|
||||
buf == buf_cloned
|
||||
), f"{buf=} {buf_cloned=} {kwargs['loc'].to_list()=}"
|
||||
return
|
||||
|
||||
cls.triton(*args, **kwargs, buf=buf)
|
||||
|
||||
@classmethod
|
||||
def vanilla(cls, pool, buf, loc, index_k, index_k_scale):
|
||||
SetK.execute(pool=pool, buf=buf, loc=loc, index_k=index_k)
|
||||
SetS.execute(pool=pool, buf=buf, loc=loc, index_k_scale=index_k_scale)
|
||||
|
||||
@classmethod
|
||||
def triton(cls, pool, buf, loc, index_k, index_k_scale):
|
||||
loc = loc.to(torch.int64)
|
||||
|
||||
@@ -252,8 +252,7 @@ def pre_permute_standard_to_triton(
|
||||
running_state: dict,
|
||||
) -> TritonRunnerInput:
|
||||
|
||||
# NOTE: this is dead code as a fused func for standard format is registered.
|
||||
# This is left here for testing and examples.
|
||||
# Registered fallback for format-conversion tests and examples.
|
||||
|
||||
from sglang.srt.layers.moe.moe_runner.triton_utils.fused_moe import (
|
||||
_prepare_fused_moe_run,
|
||||
@@ -309,8 +308,7 @@ def post_permute_triton_to_standard(
|
||||
running_state: dict,
|
||||
) -> StandardCombineInput:
|
||||
|
||||
# NOTE: this is dead code as a fused func for standard format is registered.
|
||||
# This is left here for testing and examples.
|
||||
# Registered fallback for format-conversion tests and examples.
|
||||
|
||||
from sglang.srt.layers.moe.token_dispatcher.standard import StandardCombineInput
|
||||
|
||||
|
||||
@@ -195,7 +195,7 @@ if _use_aiter:
|
||||
if _is_musa:
|
||||
try:
|
||||
from mate import moe_fused_gate
|
||||
except ImportError as e:
|
||||
except ImportError:
|
||||
raise ImportError("mate is required for the biased grouped topk.")
|
||||
|
||||
from sglang.srt.hardware_backend.musa.kernels.topk import topk_sigmoid, topk_softmax
|
||||
|
||||
@@ -94,11 +94,7 @@ def hash_tiles32_kernel_blocked(
|
||||
h1 ^= nbytes
|
||||
h2 ^= nbytes
|
||||
h1 = _fmix32(h1, C1=FM_C1, C2=FM_C2)
|
||||
h2 = (
|
||||
_fmix32(h2, C1=FMIX32_C1, C2=FMIX32_C2)
|
||||
if False
|
||||
else _fmix32(h2, C1=FM_C1, C2=FM_C2)
|
||||
)
|
||||
h2 = _fmix32(h2, C1=FM_C1, C2=FM_C2)
|
||||
|
||||
out = (h1.to(tl.uint64) << 32) | h2.to(tl.uint64)
|
||||
tl.store(out_ptr + pid, out)
|
||||
|
||||
@@ -9,7 +9,7 @@ import torch
|
||||
|
||||
try:
|
||||
from aiter.ops.triton.quant import dynamic_mxfp4_quant
|
||||
except ImportError as err:
|
||||
except ImportError:
|
||||
|
||||
def raise_aiter_import_error(*args, **kwargs):
|
||||
raise ImportError(
|
||||
|
||||
@@ -68,7 +68,7 @@ def init_feature_buffer(device):
|
||||
num_elements, dtype=torch.float32, device=device
|
||||
)
|
||||
logger.info(f"Preallocated {size_mb}MB GPU buffer")
|
||||
except RuntimeError as e:
|
||||
except RuntimeError:
|
||||
_GPU_FEATURE_BUFFER = None
|
||||
|
||||
|
||||
@@ -160,7 +160,7 @@ class TransportProxyTensor(torch.Tensor):
|
||||
"storage_offset": self.storage_offset(),
|
||||
}
|
||||
state["tensor_data"] = None
|
||||
except Exception as e:
|
||||
except Exception:
|
||||
# Failed to get CUDA IPC handle (possibly tp). Falling back to default transport.
|
||||
state["metadata"]["transport_mode"] = "default"
|
||||
state["tensor_data"] = self.as_subclass(torch.Tensor)
|
||||
|
||||
@@ -1064,7 +1064,7 @@ class ModelRunner(ModelRunnerKVCacheMixin):
|
||||
def remote_instance_init_transfer_engine(self):
|
||||
try:
|
||||
from mooncake.engine import TransferEngine
|
||||
except ImportError as e:
|
||||
except ImportError:
|
||||
logger.warning(
|
||||
"Please install mooncake for using remote instance transfer engine: pip install mooncake-transfer-engine"
|
||||
)
|
||||
|
||||
@@ -531,22 +531,15 @@ class HunYuanModel(nn.Module):
|
||||
hidden_states = self.get_input_embeddings(input_ids)
|
||||
residual = None
|
||||
|
||||
prev_kv_states = None
|
||||
for i in range(len(self.layers)):
|
||||
layer = self.layers[i]
|
||||
hidden_states, residual, kv_states = layer(
|
||||
for layer in self.layers:
|
||||
hidden_states, residual, _ = layer(
|
||||
positions,
|
||||
hidden_states,
|
||||
forward_batch,
|
||||
residual,
|
||||
prev_kv_states,
|
||||
None,
|
||||
)
|
||||
|
||||
if False: # (i - self.start_layer) % cla_factor == 0:
|
||||
prev_kv_states = kv_states
|
||||
else:
|
||||
prev_kv_states = None
|
||||
|
||||
hidden_states, _ = self.norm(hidden_states, residual)
|
||||
return hidden_states
|
||||
|
||||
|
||||
@@ -213,7 +213,7 @@ def process_anyres_image(image, processor, grid_pinpoints):
|
||||
if isinstance(grid_pinpoints, str) and "x" in grid_pinpoints:
|
||||
try:
|
||||
patch_size = processor.size[0]
|
||||
except Exception as e:
|
||||
except Exception:
|
||||
patch_size = processor.size["shortest_edge"]
|
||||
assert patch_size in [
|
||||
224,
|
||||
|
||||
@@ -1857,7 +1857,7 @@ def get_npu_memory_capacity():
|
||||
return envs.SGLANG_ZBAL_LOCAL_MEM_SIZE.get() # unit: MB
|
||||
else:
|
||||
return torch.npu.mem_get_info()[1] // 1024 // 1024 # unit: MB
|
||||
except ImportError as e:
|
||||
except ImportError:
|
||||
raise ImportError("torch_npu is required when run on npu device.")
|
||||
|
||||
|
||||
@@ -2210,7 +2210,7 @@ def get_compiler_backend(mode=None) -> str:
|
||||
import torchair
|
||||
import torchair.ge_concrete_graph.ge_converter.experimental.patch_for_hcom_allreduce
|
||||
from torchair.configs.compiler_config import CompilerConfig
|
||||
except ImportError as e:
|
||||
except ImportError:
|
||||
raise ImportError(
|
||||
"NPU detected, but torchair package is not installed. "
|
||||
"Please install torchair for torch.compile support on NPU."
|
||||
|
||||
@@ -364,7 +364,7 @@ class CudaIpcTensorTransportProxy:
|
||||
"recons_dtype": info_data.dtype,
|
||||
}
|
||||
state["tensor_data"] = None
|
||||
except Exception as e:
|
||||
except Exception:
|
||||
# Failed to get CUDA IPC handle (possibly tp). Falling back to default transport.
|
||||
state["ipc_extra"] = None
|
||||
state["tensor_data"] = data
|
||||
|
||||
@@ -368,11 +368,11 @@ def rpd_to_chrome_trace(
|
||||
if row[0] == "0": # Frame start
|
||||
if key not in stacks:
|
||||
stacks[key] = []
|
||||
stack = stacks[key].append((row[1], row[4]))
|
||||
stacks[key].append((row[1], row[4]))
|
||||
# print(f"0: new api frame: pid_tid={key} -> stack={stacks}")
|
||||
|
||||
elif row[0] == "1": # Frame end
|
||||
completed = stacks[key].pop()
|
||||
stacks[key].pop()
|
||||
# print(f"1: end api frame: pid_tid={key} -> stack={stacks}")
|
||||
|
||||
elif row[0] == "2": # API + Op
|
||||
@@ -400,8 +400,6 @@ def rpd_to_chrome_trace(
|
||||
or abs(gpuFrame.start - row[8]) < 200
|
||||
)
|
||||
):
|
||||
# if gpuFrame.id == frame[0] and gpuFrame.name == frame[1]: # Another op under the same frame -> union them
|
||||
# if False: # Turn off frame joining
|
||||
if row[7] < gpuFrame.start:
|
||||
gpuFrame.start = row[7]
|
||||
if row[8] > gpuFrame.end:
|
||||
|
||||
@@ -779,7 +779,7 @@ def wait_server_ready(url, timeout=LOCAL_TIMEOUT):
|
||||
logger.info(
|
||||
f"Server {url} returned status code: {response.status_code}"
|
||||
)
|
||||
except Exception as e:
|
||||
except Exception:
|
||||
# logger.error(f"Server {url} request error: {e}, retrying...")
|
||||
pass
|
||||
|
||||
|
||||
@@ -1104,7 +1104,7 @@ def _seed_c4_if_needed(
|
||||
fixture: DSV4AttentionFixture, *, num_entries: int = _DSV4_EXTRA_ENTRIES
|
||||
) -> None:
|
||||
"""For compress_ratio=4, seed the C4 metadata the exercised path consumes
|
||||
(the C4Indexer would normally populate it; the smoke fixture skips the
|
||||
(the C4Indexer would normally populate it; the compact fixture skips the
|
||||
indexer): `c4_sparse_page_indices` for the dense extend path,
|
||||
`c4_sparse_raw_indices` for sparse prefill. No-op for other compress_ratios.
|
||||
"""
|
||||
@@ -1383,7 +1383,7 @@ def _seed_c4_sparse_indices(
|
||||
) -> None:
|
||||
"""For compress_ratio=4 the production `init_flashmla_related` initializes
|
||||
`c4_sparse_page_indices` to all `-1` (the C4Indexer fills it in later).
|
||||
Since the smoke fixture does not run the indexer, the C4 path attends to
|
||||
Since the compact fixture does not run the indexer, the C4 path attends to
|
||||
zero extra entries unless we seed the indices ourselves. Seed each query
|
||||
row to point to `[0, 1, ..., num_entries - 1]` so the backend reads the
|
||||
same `num_entries` C4 K's that the reference also reads, exercising the
|
||||
@@ -1614,7 +1614,7 @@ def run_dsv4_compress_attention_case(
|
||||
assert case.compress_ratio in (
|
||||
4,
|
||||
128,
|
||||
), f"smoke runner requires compress_ratio in (4, 128); got {case.compress_ratio}"
|
||||
), f"DSV4 compact runner requires compress_ratio in (4, 128); got {case.compress_ratio}"
|
||||
if sparse_prefill:
|
||||
assert (
|
||||
case.forward_mode.is_extend_without_speculative()
|
||||
|
||||
@@ -2,7 +2,7 @@
|
||||
|
||||
Probes that catch the model producing wrong output: weight load
|
||||
failure, sampling path bugs, KV / attention corruption, and cuda graph
|
||||
edge cases. Single-prompt smoke only -- dataset-driven accuracy gates
|
||||
edge cases. Single-prompt coverage only -- dataset-driven accuracy gates
|
||||
belong to the consuming test class, not this kit.
|
||||
|
||||
Mix into any ``CustomTestCase`` subclass that exposes ``self.base_url``
|
||||
|
||||
@@ -175,7 +175,7 @@ class MGSMEval(Eval):
|
||||
]
|
||||
try:
|
||||
response_text = sampler(prompt_messages)
|
||||
except Exception as e:
|
||||
except Exception:
|
||||
response_text = ""
|
||||
|
||||
answer_prefix = LANG_TO_ANSWER_PREFIX[language]
|
||||
|
||||
Reference in New Issue
Block a user