guberti commented on code in PR #13752:
URL: https://github.com/apache/tvm/pull/13752#discussion_r1101019261
##########
python/tvm/topi/arm_cpu/qnn.py:
##########
@@ -368,3 +389,139 @@ def kernel_ptr(buffer, c, offset=0):
def schedule_qnn_depthwise_conv2d(_attrs, _outs, _target):
"""Schedule function for qnn.depthwise_conv2d."""
return None
+
+
+def _make_unrolled_conv2d_primfunc(
+ output_dimensions: Tuple[int, int, int],
+ buffer_shapes: Tuple[Tuple, Tuple, Tuple, Tuple, Tuple],
+ function_names: Dict[Tuple, str],
+ function_code: str,
+ ptr_gens: Tuple[Callable, Callable],
+ output_layout="NHWC",
+):
+ out_height, out_width, out_channels = output_dimensions
+ data_shape, kernel_shape, bias_shape, scale_shape, output_shape =
buffer_shapes
+ data_ptr, kernel_ptr = ptr_gens
+
+ def output_ptr(output, y, c):
+ if output_layout == "NHWC":
+ return _make_tscript_ptr(output, y * const(out_width *
out_channels) + c, 1)
+ elif output_layout == "NCHW":
+ return _make_tscript_ptr(
+ output, c * const(out_height * out_width) + y *
const(out_width), 1
+ )
+ else:
+ raise TVMError(f"Unsupported out_layout '{output_layout}'!")
+
+ def make_row_call(buffers, c_var, y, c):
+ output, data, kernel, bias, scale = buffers
+ return _make_tscript_call(
+ function_names[(y + c) % 2, c % 2, 0],
+ output_ptr(output, y, c_var + c),
+ data_ptr(data, y, c_var + c, offset=(y + c) % 2),
+ kernel_ptr(kernel, c_var + c, offset=c),
+ _bias_ptr(bias, c_var + c),
+ _scale_ptr(scale, c_var + c),
+ )
+
+ @T.prim_func
+ def biased_quantized_conv2d(
+ data_handle: T.handle,
+ kernel_handle: T.handle,
+ bias_handle: T.handle,
+ scale_handle: T.handle,
+ output_handle: T.handle,
+ ) -> None:
+ # Same setup is used as in _make_conv2d_primfunc
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ data = T.match_buffer(data_handle, data_shape, dtype="int16")
+ kernel = T.match_buffer(kernel_handle, kernel_shape, dtype="int16")
+ bias = T.match_buffer(bias_handle, bias_shape, dtype="int32")
+ scale = T.match_buffer(scale_handle, scale_shape)
+ output = T.match_buffer(output_handle, output_shape, dtype="int16")
+
+ # pylint: disable=unused-variable
+ output[0, 0, 0, 0] = 0
+ __1 = data[0, 0, 0, 0]
+ __2 = kernel[0, 0, 0, 0]
+ __3 = bias[0, 0, 0, 0]
+ __4 = scale[0]
+ # pylint: enable=unused-variable
+
+ for c_ax in T.grid(out_channels // 2):
+ with T.block("conv2ds"):
+ T.block_attr({"pragma_import_c": function_code})
+ c = T.axis.remap("S", [c_ax]) * 2
+
+ # TODO how can I programatically make the right number of
Review Comment:
Fixed!
--
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]
