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

1import logging 

2 

3import triton 

4from _kunlunxin.utils.codegen_config_utils import CodeGenConfig 

5 

6from ..utils.pointwise_dynamic import pointwise_dynamic 

7 

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

9 

10config_ = CodeGenConfig( 

11 512, 

12 (65536, 65536, 65536), 

13 32, 

14 True, 

15 prefer_1d_tile=True, 

16 isCloseMemoryAsync=False, 

17 kunlunAutoGrid=True, 

18 unroll_num=8, 

19) 

20 

21 

22@pointwise_dynamic(promotion_methods=[(0, 1, "DEFAULT")], config=config_) 

23@triton.jit 

24def bitwise_and_func(x, y): 

25 return x & y 

26 

27 

28def bitwise_and_tensor(A, B): 

29 logger.debug("GEMS BITWISE AND") 

30 return bitwise_and_func(A, B) 

31 

32 

33def bitwise_and_tensor_(A, B): 

34 logger.debug("GEMS BITWISE AND_") 

35 return bitwise_and_func(A, B, out0=A) 

36 

37 

38@pointwise_dynamic(is_tensor=[True, False], promotion_methods=[(0, 1, "DEFAULT")]) 

39@triton.jit 

40def bitwise_and_func_scalar(x, y): 

41 return x & y 

42 

43 

44def bitwise_and_scalar(A, B): 

45 logger.debug("GEMS BITWISE AND SCALAR") 

46 return bitwise_and_func_scalar(A, B) 

47 

48 

49def bitwise_and_scalar_(A, B): 

50 logger.debug("GEMS BITWISE AND_ SCALAR") 

51 return bitwise_and_func_scalar(A, B, out0=A) 

52 

53 

54def bitwise_and_scalar_tensor(A, B): 

55 logger.debug("GEMS BITWISE AND SCALAR TENSOR") 

56 return bitwise_and_func_scalar(B, A)