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

1import logging 

2 

3import triton 

4import triton.language as tl 

5 

6from ..utils.codegen_config_utils import CodeGenConfig 

7from ..utils.pointwise_dynamic import pointwise_dynamic 

8 

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

10 

11# Custom config with memory async enabled for better memory throughput 

12_abs_config = CodeGenConfig( 

13 max_tile_size=512, 

14 max_grid_size=(65536, 65536, 65536), 

15 max_num_warps_per_cta=32, 

16 prefer_block_pointer=True, 

17 prefer_1d_tile=True, 

18 isCloseMemoryAsync=False, # Enable memory async for better overlap 

19) 

20 

21 

22@pointwise_dynamic(promotion_methods=[(0, "COMPLEX_TO_FLOAT")], config=_abs_config) 

23@triton.jit 

24def abs_func(x): 

25 return tl.abs(x) 

26 

27 

28def abs(A): 

29 logger.debug("GEMS_KUNLUNXIN ABS") 

30 return abs_func(A) 

31 

32 

33def abs_(A): 

34 logger.debug("GEMS_KUNLUNXIN ABS_") 

35 abs_func(A, out0=A) 

36 return A