Coverage for src/flag_gems/runtime/backend/_kunlunxin/ops/hstack.py: 0%

1import itertools 

2import logging 

3from typing import List, Tuple, Union 

4 

5import torch 

6import triton 

7 

8from flag_gems.utils.tensor_wrapper import StridedBuffer 

9 

10from ..utils.pointwise_dynamic import pointwise_dynamic 

11 

12logger = logging.getLogger("flag_gems").getChild(__name__.lstrip(".")) 

13 

14 

15@pointwise_dynamic(is_tensor=[True], promotion_methods=[(0, "DEFAULT")]) 

16@triton.jit 

17def copy_func(x): 

18 return x 

19 

20 

21def hstack( 

22 tensors: Union[Tuple[torch.Tensor, ...], List[torch.Tensor]] 

23) -> torch.Tensor: 

24 logger.debug("GEMS HSTACK") 

25 

26 if len(tensors) == 0: 

27 raise RuntimeError("hstack expected a non-empty TensorList") 

28 

29 if tensors[0].ndim == 0: 

30 tensors[0] = tensors[0].view(1) 

31 inp0_shape = tensors[0].shape 

32 out_shape = list(inp0_shape) 

33 inp_shapes = [inp0_shape] 

34 

35 if len(inp0_shape) == 1: 

36 dim = 0 

37 else: 

38 dim = 1 

39 

40 for tensor_num, tensor in enumerate(tensors[1:]): 

41 if tensor.ndim == 0: 

42 tensor = tensor.view(1) 

43 if tensor.ndim != tensors[0].ndim: 

44 raise RuntimeError( 

45 f"Tensors must have same number of dimensions: got {tensors[0].ndim} and {tensor.ndim}" 

46 ) 

47 

48 inp_shape = tensor.shape 

49 inp_shapes.append(inp_shape) 

50 

51 for i in range(len(inp_shape)): 

52 if i != dim and inp_shape[i] != inp0_shape[i]: 

53 raise RuntimeError( 

54 f"Sizes of tensors must match except in dimension {dim}. \ 

55 Expected size {inp0_shape[i]} but got size {inp_shape[i]} \ 

56 for tensor number {tensor_num + 1} in the list." 

57 ) 

58 

59 out_shape[dim] = sum(s[dim] for s in inp_shapes) 

60 

61 out0 = torch.empty(out_shape, dtype=tensors[0].dtype, device=tensors[0].device) 

62 out0_strides = out0.stride() 

63 out0_offsets = list( 

64 itertools.accumulate( 

65 [s[dim] * out0_strides[dim] for s in inp_shapes[:-1]], initial=0 

66 ) 

67 ) 

68 

69 for a, out0_offset in zip(tensors, out0_offsets): 

70 in_view = StridedBuffer(a, a.shape, a.stride()) 

71 out_view = StridedBuffer(out0, a.shape, out0.stride(), offset=out0_offset) 

72 copy_func.instantiate(a.ndim)(in_view, out0=out_view) 

73 

74 return out0