Add aten::maxunpool2d and aten::maxunpool3d

src/ATen/native/xpu/MaxUnpooling.cpp

@@ -0,0 +1,59 @@
+#include <ATen/ATen.h>
+#include <ATen/AccumulateType.h>
+#include <ATen/core/Tensor.h>
+#include <ATen/core/op_registration/adaption.h>
+#include <ATen/native/layer_norm.h>
+#include <ATen/native/xpu/sycl/MaxUnpoolingKernels.h>
+
+#ifndef AT_PER_OPERATOR_HEADERS
+#include <ATen/Functions.h>
+#include <ATen/NativeFunctions.h>
+#else
+#include <ATen/ops/empty.h>
+#endif
+
+namespace at::native {
+
+Tensor& max_unpooling2d_forward_out_xpu(
+    const Tensor& self,
+    const Tensor& indices,
+    IntArrayRef output_size,
+    Tensor& out) {
+  native::xpu::max_unpooling2d_forward_kernel(out, self, indices, output_size);
+  return out;
+}
+
+Tensor max_unpooling2d_forward_xpu(
+    const Tensor& self,
+    const Tensor& indices,
+    IntArrayRef output_size) {
+  auto out = at::empty({0}, self.options());
+  max_unpooling2d_forward_out_xpu(self, indices, output_size, out);
+  return out;
+}
+
+Tensor& max_unpooling3d_forward_out_xpu(
+    const Tensor& self,
+    const Tensor& indices,
+    IntArrayRef output_size,
+    IntArrayRef stride,
+    IntArrayRef padding,
+    Tensor& out) {
+  native::xpu::max_unpooling3d_forward_kernel(
+      out, self, indices, output_size, stride, padding);
+  return out;
+}
+
+Tensor max_unpooling3d_forward_xpu(
+    const Tensor& self,
+    const Tensor& indices,
+    IntArrayRef output_size,
+    IntArrayRef stride,
+    IntArrayRef padding) {
+  auto out = at::empty({0}, self.options());
+  max_unpooling3d_forward_out_xpu(
+      self, indices, output_size, stride, padding, out);
+  return out;
+}
+
+} // namespace at::native

src/ATen/native/xpu/XPUFallback.template

@@ -208,8 +208,6 @@ TORCH_LIBRARY_IMPL(aten, XPU, m) {
     "lu_unpack.out",
     "max_pool3d_with_indices",
     "max_pool3d_with_indices_backward",
-    "max_unpool2d",
-    "max_unpool3d",
     "multilabel_margin_loss_backward",
     "multilabel_margin_loss_forward",
     "multi_margin_loss",

src/ATen/native/xpu/sycl/KernelUtils.h

@@ -0,0 +1,13 @@
+#pragma once
+
+#include <c10/util/Exception.h>
+#include <limits>
+
+#define XPU_KERNEL_LOOP_TYPE(item, i, n, index_type)                      \
+  int64_t _i_n_d_e_x =                                                    \
+      item.get_group(0) * item.get_local_range(0) + item.get_local_id(0); \
+  for (index_type i = _i_n_d_e_x; _i_n_d_e_x < (n);                       \
+       _i_n_d_e_x += item.get_local_range(0) * item.get_group_range(0),   \
+                  i = _i_n_d_e_x)
+
+#define XPU_KERNEL_LOOP(item, i, n) XPU_KERNEL_LOOP_TYPE(item, i, n, int)