Implement average pooling

Author: parthchadha

tripy/tests/integration/test_pooling.py

@@ -31,13 +31,20 @@ class TestPooling:
         ],
     )
     @pytest.mark.parametrize("dtype", [tp.float32, tp.float16, tp.int8])
-    def test_maxpool_2d(self, kernel_dims, stride, padding, dtype):
+    @pytest.mark.parametrize("pool_type", ["max", "avg"])
+    def test_pool_2d(self, kernel_dims, stride, padding, dtype, pool_type):
         inp_tp = tp.reshape(tp.arange(64, dtype=dtype), (1, 1, 8, 8))
-        out = tp.maxpool(inp_tp, kernel_dims=kernel_dims, stride=stride, padding=padding)
-        out_torch = torch.from_dlpack(out).to("cpu")
-
         torch_padding = (padding[0][0], padding[1][0])
-        pool_torch = torch.nn.MaxPool2d(kernel_size=kernel_dims, stride=stride, padding=torch_padding)
+
+        if pool_type == "max":
+            out = tp.maxpool(inp_tp, kernel_dims=kernel_dims, stride=stride, padding=padding)
+            pool_torch = torch.nn.MaxPool2d(kernel_size=kernel_dims, stride=stride, padding=torch_padding)
+        elif pool_type == "avg":
+            pytest.skip("https://github.com/NVIDIA/TensorRT-Incubator/issues/237: Average pooling is not functional.")
+            out = tp.avgpool(inp_tp, kernel_dims=kernel_dims, stride=stride, padding=padding)
+            pool_torch = torch.nn.AvgPool2d(kernel_size=kernel_dims, stride=stride, padding=torch_padding)
+
+        out_torch = torch.from_dlpack(out).to("cpu")
         expected = pool_torch(torch.from_dlpack(inp_tp).to("cpu"))
         assert torch.allclose(expected, out_torch)
         assert expected.shape == out_torch.shape
@@ -49,13 +56,21 @@ def test_maxpool_2d(self, kernel_dims, stride, padding, dtype):
         ],
     )
     @pytest.mark.parametrize("dtype", [tp.float32, tp.float16])
-    def test_maxpool_3d(self, kernel_dims, stride, padding, dtype):
+    @pytest.mark.parametrize("pool_type", ["max", "avg"])
+    def test_pool_3d(self, kernel_dims, stride, padding, dtype, pool_type):
         inp_tp = tp.reshape(tp.arange(512, dtype=dtype), (1, 1, 8, 8, 8))
-        out = tp.maxpool(inp_tp, kernel_dims=kernel_dims, stride=stride, padding=padding)
+        torch_padding = (padding[0][0], padding[1][0], padding[2][0])
+
+        if pool_type == "max":
+            out = tp.maxpool(inp_tp, kernel_dims=kernel_dims, stride=stride, padding=padding)
+            pool_torch = torch.nn.MaxPool3d(kernel_size=kernel_dims, stride=stride, padding=torch_padding)
+        elif pool_type == "avg":
+            pytest.skip("https://github.com/NVIDIA/TensorRT-Incubator/issues/237: Average pooling is not functional.")
+            out = tp.avgpool(inp_tp, kernel_dims=kernel_dims, stride=stride, padding=padding)
+            pool_torch = torch.nn.AvgPool3d(kernel_size=kernel_dims, stride=stride, padding=torch_padding)
+
         out_torch = torch.from_dlpack(out).to("cpu")
 
-        torch_padding = (padding[0][0], padding[1][0], padding[2][0])
-        pool_torch = torch.nn.MaxPool3d(kernel_size=kernel_dims, stride=stride, padding=torch_padding)
         expected = pool_torch(torch.from_dlpack(inp_tp).to("cpu"))
         assert torch.allclose(expected, out_torch)
         assert expected.shape == out_torch.shape

tripy/tripy/flat_ir/ops/reduce_window.py

@@ -36,6 +36,8 @@ class ReduceWindowOp(BaseFlatIROp):
     def _get_reduce_func(self):
         if self.reduce_mode == "max":
             return stablehlo.MaxOp
+        elif self.reduce_mode == "avg":
+            return stablehlo.AddOp
         else:
             raise NotImplementedError()
 
@@ -71,6 +73,21 @@ def to_mlir(self, operands):
         reduce_func = self._get_reduce_func()
         with ir.InsertionPoint(reduce_block):
             out = reduce_func(*reduce_block.arguments)
+            if self.reduce_mode == "avg":
+                # Calculate the number of elements in the window
+                window_elements = 1
+                for dim in self.window_dims:
+                    window_elements *= dim
+
+                # Create a dense elements attribute for the window size
+                window_size_attr = ir.DenseElementsAttr.get_splat(
+                    ir.RankedTensorType.get([], mlir_utils.get_mlir_dtype(input_dtype)),
+                    ir.FloatAttr.get(mlir_utils.get_mlir_dtype(input_dtype), float(window_elements)),
+                )
+                window_size_const = stablehlo.ConstantOp(window_size_attr)
+
+                # Divide the sum by the window size
+                out = stablehlo.DivOp(out, window_size_const)
             stablehlo.ReturnOp([out])
 
         return [reduce]

tripy/tripy/frontend/trace/ops/pooling.py

@@ -106,7 +106,7 @@ def maxpool(
     Args:
         input: The input tensor.
         kernel_dims: The spatial shape of the pooling window. Only 2-D or 3-D ``kernel_dims`` are supported.
-            If the input has ``int8`` datatype, ``kernel_dims`` can only be 2-D.
+            If the input has :class:`int8` datatype, ``kernel_dims`` can only be 2-D.
         stride: A sequence of length :math:`M` indicating the stride of pooling across each spatial dimension,
             where :math:`M` is the number of spatial dimensions, i.e. :math:`M = \text{rank(input)} - 2`.
             Defaults to all 1.
@@ -139,3 +139,64 @@ def maxpool(
     padding = utils.default(padding, [(0, 0)] * spatial_dims)
 
     return Pooling.build([input], Pooling.Kind.MAX, kernel_dims, stride, padding)
+
+
+@export.public_api(document_under="operations/functions")
+@constraints.dtype_info(
+    dtype_variables={
+        "T1": ["float32", "float16", "int8"],
+    },
+    dtype_constraints={"input": "T1", constraints.RETURN_VALUE: "T1"},
+)
+def avgpool(
+    input: "tripy.Tensor",
+    kernel_dims: Sequence[int],
+    stride: Sequence[int] = None,
+    padding: Sequence[Tuple[int]] = None,
+) -> "tripy.Tensor":
+    r"""
+    Applies an average pooling over the input tensor.
+
+    The output's non-spatial dimensions are the same as input. For each input spatial dimension
+    :math:`D_{i}`, the corresponding output dimension will be:
+
+    .. math::
+        D_{out_i} = \left\lfloor\frac{D_{i} + \text{padding_before[i]} + \text{padding_after[i]} -
+                \text{kernel_dims[i]}}{\text{stride[i]}} + 1\right\rfloor
+
+    Args:
+        input: The input tensor.
+        kernel_dims: The spatial shape of the pooling window. Only 2-D or 3-D ``kernel_dims`` are supported.
+            If the input has :class:`int8` datatype, ``kernel_dims`` can only be 2-D.
+        stride: A sequence of length :math:`M` indicating the stride of pooling across each spatial dimension,
+            where :math:`M` is the number of spatial dimensions, i.e. :math:`M = \text{rank(input)} - 2`.
+            Defaults to all 1.
+        padding: A sequence of pairs of integers of length :math:`M` indicating the zero padding
+            to apply to the input along each spatial dimension before and after the dimension respectively,
+            where :math:`M` is the number of spatial dimensions, i.e. :math:`M = \text{rank(input)} - 2`.
+            Defaults to all 0.
+
+    Returns:
+        The result tensor after the pooling operation.
+
+    .. code-block:: python
+        :linenos:
+        :caption: Example
+
+        input = tp.reshape(tp.arange(16, dtype=tp.float32), (1, 1, 4, 4))
+        output = tp.avgpool(input, kernel_dims=(2, 2))
+
+        pool_torch = torch.nn.AvgPool2d((2, 2), stride=1) # doc: omit
+        expected = pool_torch(torch.from_dlpack(input).to("cpu")) # doc: omit
+
+        assert torch.allclose(torch.from_dlpack(output).to("cpu"), expected)
+    """
+    spatial_dims = len(kernel_dims)
+    if spatial_dims != 2 and spatial_dims != 3:
+        raise_error("Unsupported kernel_dims, must be 2D or 3D.", [f"Got kernel_dims={kernel_dims}"])
+
+    op_utils.check_conv_pooling_args(kernel_dims, stride, padding)
+    stride = utils.default(stride, [1] * spatial_dims)
+    padding = utils.default(padding, [(0, 0)] * spatial_dims)
+
+    return Pooling.build([input], Pooling.Kind.AVG, kernel_dims, stride, padding)