(
x,
dx,
dy,
T: tl.constexpr,
D: tl.constexpr,
BT: tl.constexpr
)
| 106 | ) |
| 107 | @triton.jit |
| 108 | def logsigmoid_bwd_kernel( |
| 109 | x, |
| 110 | dx, |
| 111 | dy, |
| 112 | T: tl.constexpr, |
| 113 | D: tl.constexpr, |
| 114 | BT: tl.constexpr |
| 115 | ): |
| 116 | i = tl.program_id(0) |
| 117 | o_i = i * BT + tl.arange(0, BT) |
| 118 | |
| 119 | p_x = x + o_i |
| 120 | p_dx = dx + o_i |
| 121 | p_dy = dy + o_i |
| 122 | mask = o_i < T |
| 123 | |
| 124 | # [D,] |
| 125 | b_x = tl.load(p_x, mask=mask, other=0.).to(tl.float32) |
| 126 | b_dy = tl.load(p_dy, mask=mask, other=0.).to(tl.float32) |
| 127 | b_dx = b_dy * (1. - tl.sigmoid(b_x)) |
| 128 | tl.store(p_dx, b_dx.to(p_dx.dtype.element_ty), mask=mask) |
| 129 | |
| 130 | |
| 131 | class LogSigmoidFunction(torch.autograd.Function): |
nothing calls this directly
no outgoing calls
no test coverage detected