[INTERPRETER] Support unary ops #3279
Conversation
| @@ -244,7 +244,7 @@ def kernel(Z, X, SIZE: tl.constexpr): | |||
| z_ref = eval(expr if numpy_expr is None else numpy_expr) | |||
| # triton result | |||
| x_tri = to_triton(x, device=device, dst_type=dtype_x) | |||
| z_tri = to_triton(np.empty_like(z_ref), device=device, dst_type=dtype_x) | |||
| z_tri = to_triton(np.empty_like(x), device=device, dst_type=dtype_x) | |||
There was a problem hiding this comment.
FYI. There was a memory leak in our tests
There was a problem hiding this comment.
For my info, how does this leak memory?
There was a problem hiding this comment.
Let z_ref be a scalar, np.empty_like is also a scalar, and so is z_tri. to_triton doesn't splat a scalar, but only reinterpret a value to the given dst_type. However, the output of the kernel on the next line will be an array of 128 elements.
I'm wondering whether we have similar problems at other places.
There was a problem hiding this comment.
Ah, it's an out-of-bounds memory access (I usually reserve "memory leak" for memory that's allocated and never freed, but maybe you use a different definition). btw .contiguous() is another way to fix this.
There was a problem hiding this comment.
Probably to_triton should just call contiguous().
No description provided.