Implement average pooling

parthchadha · parthchadha · commit 4b7ca047abd0 · 2024-10-01T08:42:47.000-07:00
diff --git a/tripy/tests/integration/test_pooling.py b/tripy/tests/integration/test_pooling.py
@@ -31,13 +31,26 @@ 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":
+            if torch_padding != (0, 0):
+                pytest.skip(
+                    "https://github.com/NVIDIA/TensorRT-Incubator/issues/241: Tripy average pool is incorrect when padding != 0."
+                )
+
+            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")
+        if pool_type == "avg" and dtype == tp.int8:
+            pytest.skip("Torch average pool is not implemented for int8")
         expected = pool_torch(torch.from_dlpack(inp_tp).to("cpu"))
         assert torch.allclose(expected, out_torch)
         assert expected.shape == out_torch.shape
@@ -49,13 +62,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
diff --git a/tripy/tripy/flat_ir/ops/reduce_window.py b/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()
 
diff --git a/tripy/tripy/frontend/trace/ops/pooling.py b/tripy/tripy/frontend/trace/ops/pooling.py
@@ -49,7 +49,7 @@ def infer_dtypes(self):
         self.outputs[0].dtype = self.inputs[0].dtype
 
     def to_flat_ir(self, inputs, outputs):
-        from tripy.flat_ir.ops import ConstantOp, ReduceWindowOp
+        from tripy.flat_ir.ops import ConstantOp, DivideOp, ReduceWindowOp
         from tripy.flat_ir.tensor import FlatIRTensor
 
         init_value = 0
@@ -70,14 +70,74 @@ def to_flat_ir(self, inputs, outputs):
         window_strides = [1] * extra_dims + list(self.stride)
         padding = [(0, 0)] * extra_dims + list(self.padding)
 
-        ReduceWindowOp.build(
-            [inputs[0], init_const],
-            outputs,
-            reduce_mode=self.kind.op,
-            window_dims=window_dims,
-            window_strides=window_strides,
-            padding=padding,
-        )
+        if self.kind.op == "max":
+            ReduceWindowOp.build(
+                [inputs[0], init_const],
+                outputs,
+                reduce_mode=self.kind.op,
+                window_dims=window_dims,
+                window_strides=window_strides,
+                padding=padding,
+            )
+        elif self.kind.op == "avg":
+
+            reduce_out = FlatIRTensor.build(
+                rank=outputs[0].rank,
+                dtype=outputs[0].dtype,
+                device=outputs[0].device,
+                reason_details=[f"create the output of reduce `{self.kind.op}` operation."],
+            )
+
+            ReduceWindowOp.build(
+                [inputs[0], init_const],
+                [reduce_out],
+                reduce_mode=self.kind.op,
+                window_dims=window_dims,
+                window_strides=window_strides,
+                padding=padding,
+            )
+
+            window_elements = 1
+            for dim in window_dims:
+                window_elements *= dim
+
+            # window_elements = compute_window_elements(self.kernel_dims, self.padding)
+            init_const = FlatIRTensor.build(
+                shape=(),
+                rank=0,
+                dtype=outputs[0].dtype,
+                device=outputs[0].device,
+                reason_details=[
+                    f"create the constant value tensor (containing {window_elements}) for the divisor of average pool operation."
+                ],
+            )
+            ConstantOp.build([], [init_const], data=window_elements)
+            with FlatIRTensor.context(
+                [f"expand the rank of constant tensor which is the divisor of average pool operation."]
+            ):
+                init_const = op_utils.expand_rank_of_tensor(init_const, inputs[0].rank)
+
+            with FlatIRTensor.context([f"broadcast the inputs of division operation."]):
+                shape_of_input0 = op_utils.get_shape_of_tensor(reduce_out)
+                shape_of_input1 = op_utils.get_shape_of_tensor(init_const)
+
+                # Compute element-wise max of input shapes to get the desired output shape.
+                output_shape_tensor = op_utils.compute_shape_of_broadcast(
+                    shape_of_input0,
+                    shape_of_input1,
+                    inputs[0].rank,
+                    shape1_name=f"the shape of the first input {shape_of_input0}",
+                    shape2_name=f"the shape of the second input {shape_of_input1}",
+                )
+
+                init_const = op_utils.insert_broadcast(
+                    init_const,
+                    out_rank=inputs[0].rank,
+                    shape_of_target_tensor=output_shape_tensor,
+                    tensor_details=f"left operand",
+                )
+
+            DivideOp.build([reduce_out, init_const], outputs)
 
 
 @export.public_api(document_under="operations/functions")
@@ -106,7 +166,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 +199,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"],
+    },
+    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)
