Coverage for src/flag_gems/runtime/backend/_cambricon/ops/relu.py: 0%

1import logging 

2 

3import triton 

4import triton.language as tl 

5 

6from ..utils.pointwise_dynamic import pointwise_dynamic 

7 

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

9 

10 

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

12@triton.jit 

13def relu_forward(x, inplace): 

14 return tl.where(x > 0, x, 0) 

15 

16 

17@pointwise_dynamic(promotion_methods=[(0, "DEFAULT")]) 

18@triton.jit 

19def relu_backward(x, dy): 

20 return tl.where(x > 0, dy, 0) 

21 

22 

23def relu(self): 

24 logger.debug("GEMS_CAMBRICON RELU FORWARD") 

25 output = relu_forward(self, False) 

26 return output 

27 

28 

29def relu_(A): 

30 logger.debug("GEMS_CAMBRICON RELU_ FORWARD") 

31 out = relu_forward(A, True, out0=A) 

32 return out