delldu commented on a change in pull request #8443:
URL: https://github.com/apache/tvm/pull/8443#discussion_r670290905
##########
File path: include/tvm/topi/nn.h
##########
@@ -690,6 +690,130 @@ inline Tensor nll_loss(const Tensor& predictions, const
Tensor& targets, const T
return T;
}
}
+
+/*!
+ * \brief Creates an operation that performs im2col with an NCHW-layout
+ *
+ * \param data The 4-D input tensor
+ * \param kernel_size A static tuple for kernel size, such as (3,3)
+ * \param dilation A static tuple for dilation, default is (1,1)
+ * \param padding A static tuple for padding, padding value is zero
+ * \param stride A static tuple for strides, default is (1,1)
+ * \param name The name of the operation
+ * \param tag The tag to mark the operation
+ *
+ * \return A Tensor whose op member is the im2col operation (NCHW layout)
+
+ Pseudo code:
+ input_b, input_c, input_h, input_w = data_shape
+ dilation_h, dilation_w = dilation
+ padding_h, padding_w = padding
+
+ dilated_kernel_h = (kernel_h - 1) * dilation_h + 1
+ dilated_kernel_w = (kernel_w - 1) * dilation_w + 1
+
+ output_h = (input_h + 2 * padding_h - dilated_kernel_h)//stride_h + 1
+ output_w = (input_w + 2 * padding_w - dilated_kernel_w)//stride_w + 1
+
+ h_offset = (k // input_w) % kernel_h
+ w_offset = k % kernel_w
+
+ im2col_data = te.compute(
+ (N, K, L),
+ lambda n, k, l: data[
+ n,
+ k / kernel_h / kernel_w,
+ stride_h * (l / output_w) + dilation_h * h_offset - padding_h,
+ stride_w * (l % output_w) + dilation_w * w_offset - padding_w,
+ ],
+ name="im2col_data",
+ )
+*/
+inline tvm::te::Tensor im2col(const tvm::te::Tensor& data,
+ const tvm::Array<tvm::PrimExpr>& kernel_size,
+ const tvm::Array<tvm::PrimExpr>& dilation,
+ const tvm::Array<tvm::PrimExpr>& padding,
+ const tvm::Array<tvm::PrimExpr>& stride,
+ std::string name = "T_im2col", std::string tag =
kElementWise) {
+ ICHECK_EQ(4, data->shape.size());
+ ICHECK_EQ(2, kernel_size.size());
+ ICHECK_EQ(2, dilation.size());
+ ICHECK_EQ(2, padding.size());
+ ICHECK_EQ(2, stride.size());
+
+ auto input_b = data->shape[0];
+ auto input_c = data->shape[1];
+ auto input_h = data->shape[2];
+ auto input_w = data->shape[3];
+
+ auto kernel_h = tvm::cast(tvm::DataType::Int(32), kernel_size[0]);
+ auto kernel_w = tvm::cast(tvm::DataType::Int(32), kernel_size[1]);
+
+ auto dilation_h = tvm::cast(tvm::DataType::Int(32), dilation[0]);
+ auto dilation_w = tvm::cast(tvm::DataType::Int(32), dilation[1]);
+
+ auto padding_h = tvm::cast(tvm::DataType::Int(32), padding[0]);
+ auto padding_w = tvm::cast(tvm::DataType::Int(32), padding[1]);
+
+ auto stride_h = tvm::cast(tvm::DataType::Int(32), stride[0]);
+ auto stride_w = tvm::cast(tvm::DataType::Int(32), stride[1]);
Review comment:
Please reference on same file about 155~245 lines.
```
tvm::te::Tensor pad(const tvm::te::Tensor& t, const
tvm::Array<tvm::PrimExpr>& pad_before,
tvm::Array<tvm::PrimExpr> pad_after =
tvm::Array<tvm::PrimExpr>(),
PrimExpr pad_value = PrimExpr(), std::string name
= "T_pad",
std::string tag = kElementWise, std::string
pad_mode = "constant",
const Array<PrimExpr>* dyn_output_shape = nullptr)
```
with same function tvm::cast to parsing parameters.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
