|
| 1 | +import unittest |
1 | 2 | from typing import Tuple |
2 | 3 |
|
3 | 4 | import torch |
|
11 | 12 | from ..conversion.harness import DispatchTestCase |
12 | 13 |
|
13 | 14 |
|
14 | | -@unittest.skipIf( |
15 | | - torch_tensorrt.ENABLED_FEATURES.tensorrt_rtx, |
16 | | - "TensorRT RTX does not support plugins", |
17 | | -) |
18 | | -class TestAutomaticPlugin(DispatchTestCase): |
| 15 | +@triton.jit |
| 16 | +def elementwise_scale_mul_kernel(X, Y, Z, a, b, BLOCK_SIZE: tl.constexpr): |
| 17 | + pid = tl.program_id(0) |
| 18 | + # Compute the range of elements that this thread block will work on |
| 19 | + block_start = pid * BLOCK_SIZE |
| 20 | + # Range of indices this thread will handle |
| 21 | + offsets = block_start + tl.arange(0, BLOCK_SIZE) |
| 22 | + # Load elements from the X and Y tensors |
| 23 | + x_vals = tl.load(X + offsets) |
| 24 | + y_vals = tl.load(Y + offsets) |
| 25 | + # Perform the element-wise multiplication |
| 26 | + z_vals = x_vals * y_vals * a + b |
| 27 | + # Store the result in Z |
| 28 | + tl.store(Z + offsets, z_vals) |
| 29 | + |
| 30 | + |
| 31 | +@torch.library.custom_op("torchtrt_ex::elementwise_scale_mul", mutates_args=()) # type: ignore[misc] |
| 32 | +def elementwise_scale_mul( |
| 33 | + X: torch.Tensor, Y: torch.Tensor, b: float = 0.2, a: int = 2 |
| 34 | +) -> torch.Tensor: |
| 35 | + # Ensure the tensors are on the GPU |
| 36 | + assert X.is_cuda and Y.is_cuda, "Tensors must be on CUDA device." |
| 37 | + assert X.shape == Y.shape, "Tensors must have the same shape." |
| 38 | + |
| 39 | + # Create output tensor |
| 40 | + Z = torch.empty_like(X) |
| 41 | + |
| 42 | + # Define block size |
| 43 | + BLOCK_SIZE = 1024 |
| 44 | + |
| 45 | + # Grid of programs |
| 46 | + grid = lambda meta: (X.numel() // meta["BLOCK_SIZE"],) |
| 47 | + |
| 48 | + # Launch the kernel with parameters a and b |
| 49 | + elementwise_scale_mul_kernel[grid](X, Y, Z, a, b, BLOCK_SIZE=BLOCK_SIZE) |
19 | 50 |
|
20 | | - @triton.jit |
21 | | - def elementwise_scale_mul_kernel(X, Y, Z, a, b, BLOCK_SIZE: tl.constexpr): |
22 | | - pid = tl.program_id(0) |
23 | | - # Compute the range of elements that this thread block will work on |
24 | | - block_start = pid * BLOCK_SIZE |
25 | | - # Range of indices this thread will handle |
26 | | - offsets = block_start + tl.arange(0, BLOCK_SIZE) |
27 | | - # Load elements from the X and Y tensors |
28 | | - x_vals = tl.load(X + offsets) |
29 | | - y_vals = tl.load(Y + offsets) |
30 | | - # Perform the element-wise multiplication |
31 | | - z_vals = x_vals * y_vals * a + b |
32 | | - # Store the result in Z |
33 | | - tl.store(Z + offsets, z_vals) |
34 | | - |
35 | | - @torch.library.custom_op("torchtrt_ex::elementwise_scale_mul", mutates_args=()) # type: ignore[misc] |
36 | | - def elementwise_scale_mul( |
37 | | - X: torch.Tensor, Y: torch.Tensor, b: float = 0.2, a: int = 2 |
38 | | - ) -> torch.Tensor: |
39 | | - # Ensure the tensors are on the GPU |
40 | | - assert X.is_cuda and Y.is_cuda, "Tensors must be on CUDA device." |
41 | | - assert X.shape == Y.shape, "Tensors must have the same shape." |
42 | | - |
43 | | - # Create output tensor |
44 | | - Z = torch.empty_like(X) |
45 | | - |
46 | | - # Define block size |
47 | | - BLOCK_SIZE = 1024 |
48 | | - |
49 | | - # Grid of programs |
50 | | - grid = lambda meta: (X.numel() // meta["BLOCK_SIZE"],) |
51 | | - |
52 | | - # Launch the kernel with parameters a and b |
53 | | - elementwise_scale_mul_kernel[grid](X, Y, Z, a, b, BLOCK_SIZE=BLOCK_SIZE) |
54 | | - |
55 | | - return Z |
56 | | - |
57 | | - @torch.library.register_fake("torchtrt_ex::elementwise_scale_mul") |
58 | | - def _(x: torch.Tensor, y: torch.Tensor, b: float = 0.2, a: int = 2) -> torch.Tensor: |
59 | | - return x |
| 51 | + return Z |
60 | 52 |
|
| 53 | + |
| 54 | +@torch.library.register_fake("torchtrt_ex::elementwise_scale_mul") |
| 55 | +def _(x: torch.Tensor, y: torch.Tensor, b: float = 0.2, a: int = 2) -> torch.Tensor: |
| 56 | + return x |
| 57 | + |
| 58 | + |
| 59 | +if not torch_tensorrt.ENABLED_FEATURES.tensorrt_rtx: |
61 | 60 | torch_tensorrt.dynamo.conversion.plugins.custom_op( |
62 | 61 | "torchtrt_ex::elementwise_scale_mul", supports_dynamic_shapes=True |
63 | 62 | ) |
64 | 63 |
|
| 64 | + |
| 65 | +@unittest.skipIf( |
| 66 | + torch_tensorrt.ENABLED_FEATURES.tensorrt_rtx, |
| 67 | + "TensorRT RTX does not support plugins", |
| 68 | +) |
| 69 | +class TestAutomaticPlugin(DispatchTestCase): |
| 70 | + |
65 | 71 | @parameterized.expand( |
66 | 72 | [ |
67 | 73 | ((64, 64), torch.float), |
|
0 commit comments