diff --git a/vllm/distributed/kv_transfer/kv_connector/v1/hf3fs/utils/gather_scatter_helper.py b/vllm/distributed/kv_transfer/kv_connector/v1/hf3fs/utils/gather_scatter_helper.py index aebaecd40b..4ba3ecee90 100644 --- a/vllm/distributed/kv_transfer/kv_connector/v1/hf3fs/utils/gather_scatter_helper.py +++ b/vllm/distributed/kv_transfer/kv_connector/v1/hf3fs/utils/gather_scatter_helper.py @@ -4,6 +4,7 @@ import torch from vllm.logger import init_logger from vllm.triton_utils import tl, triton +from vllm.utils.torch_utils import async_tensor_h2d @triton.jit @@ -140,9 +141,9 @@ def scatter_kv_caches( assert len(kv_cache_strides) == 4 device = src_tensor.device - token_indices_tensor = torch.tensor( - token_indices, dtype=torch.int32, device="cpu" - ).to(device, non_blocking=True) + token_indices_tensor = async_tensor_h2d( + token_indices, device=device, dtype=torch.int32 + ) grid = (num_layers, num_tokens_in_block) BLOCK_SIZE = 128 @@ -196,9 +197,9 @@ def gather_kv_caches( assert len(kv_cache_strides) == 4 device = dst_tensor.device - token_indices_tensor = torch.tensor( - token_indices, dtype=torch.int32, device="cpu" - ).to(device, non_blocking=True) + token_indices_tensor = async_tensor_h2d( + token_indices, device=device, dtype=torch.int32 + ) grid = (num_layers, num_tokens_in_block) BLOCK_SIZE = 128