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

1import logging 

2 

3import triton 

4import triton.language as tl 

5 

6from flag_gems.runtime import device 

7 

8from ..utils.pointwise_dynamic import pointwise_dynamic 

9 

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

11device = device.name 

12 

13 

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

15@triton.jit 

16def maximum_kernel(X, Y): 

17 if X.dtype == tl.bfloat16: 

18 X = X.to(tl.float32) 

19 Y = Y.to(tl.float32) 

20 

21 return tl.maximum(X, Y) 

22 

23 

24def maximum(X, Y): 

25 logger.debug("GEMS MAXIMUM") 

26 assert X.device.type == device and Y.device.type == device 

27 return maximum_kernel(X, Y)