You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/06/04 08:49:14 UTC
[GitHub] [tvm] Qianshui-Jiang opened a new pull request, #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Qianshui-Jiang opened a new pull request, #11571:
URL: https://github.com/apache/tvm/pull/11571
This PR mainly about mapping oneDNN OP implementation in X86 Relay Op Strategy. we've observed that nn.dense kernel that could be dispatched to DNNL by assigning "-libs=mkldnn" and there is also conv2d kernel implemented in runtime/contrib/dnnl.
so we mapping it in X86 Relay Op Strategy and optimized the kernel implementation to let DNNL choose blocked format according to different input shape, as [performance-profiling example](https://oneapi-src.github.io/oneDNN/page_performance_profiling_cpp.html#doxid-performance-profiling-cpp) discribed in oneDNN doc.
Here is the details:
- Adjust DNNL Conv2D implementation to let it support NHWC format and automate reorder for input/weights/outputs to abtain best performace.
- Add 'target.libs=mkldnn' branch in Relay X86 OP strategy for both NCHW and NHWC Conv2D kernel.
- Add 2 test funtions for case mentioned above.
We are trying to enable more DNNL kernels including different format and datatypes this way.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Qianshui-Jiang commented on a diff in pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
Qianshui-Jiang commented on code in PR #11571:
URL: https://github.com/apache/tvm/pull/11571#discussion_r891862911
##########
python/tvm/contrib/mkldnn.py:
##########
@@ -50,3 +51,107 @@ def matmul(lhs, rhs, transa=False, transb=False, **kwargs):
name="C",
**kwargs,
)
+
+
+def dnnl_conv2d(
+ src,
+ weights,
+ stride,
+ padding,
+ dilation,
+ groups,
+ channel_last=False,
+ out_dtype="float32",
+ **kwargs,
+):
+ """Convolution operator in NCHW layout.
+
+ Parameters
+ ----------
+ src : tvm.te.Tensor
+ 4-D with shape [batch, in_channel, in_height, in_width]
+
+ weights : tvm.te.Tensor
+ 4-D with shape [num_filter, in_channel, filter_height, filter_width]
+
+ stride : int or a list/tuple of two ints
+ Stride size, or [stride_height, stride_width]
+
+ padding : int or a list/tuple of 2 or 4 ints
+ padding size, or
+ [pad_height, pad_width] for 2 ints, or
+ [pad_top, pad_left, pad_bottom, pad_right] for 4 ints
+
+ dilation: int or a list/tuple of two ints
+ dilation size, or [dilation_height, dilation_width]
+
+ groups: str
+ input data layout: NCHW or NHWC
+
+ channel_last: bool
+ chose if input/output data format is in channel_last format(NHWC) or
+ in plain format(NCHW)
+
+ out_dtype: str
+ output datatype: now only support float32
+
+ Returns
+ -------
+ Output : tvm.te.Tensor
+ 4-D with shape [batch, out_channel, out_height, out_width]
+ """
+
+ assert isinstance(stride, int) or len(stride) == 2
+ assert isinstance(dilation, int) or len(dilation) == 2
+ if isinstance(stride, int):
+ stride_h = stride_w = stride
+ else:
+ stride_h, stride_w = stride
+
+ if isinstance(dilation, int):
+ dilation_h = dilation_w = dilation
+ else:
+ dilation_h, dilation_w = dilation
+
+ if channel_last:
+ batch, in_height, in_width, _ = src.shape
+ kernel_h, kernel_w, _, num_filter = weights.shape
+ else:
+ batch, _, in_height, in_width = src.shape
+ num_filter, _, kernel_h, kernel_w = weights.shape
+
+ dilated_kernel_h = (kernel_h - 1) * dilation_h + 1
+ dilated_kernel_w = (kernel_w - 1) * dilation_w + 1
+ pad_top, pad_left, pad_down, pad_right = get_pad_tuple(
+ padding, (dilated_kernel_h, dilated_kernel_w)
+ )
+ out_channel = num_filter
+ out_height = (in_height - dilated_kernel_h + pad_top + pad_down) // stride_h + 1
+ out_width = (in_width - dilated_kernel_w + pad_left + pad_right) // stride_w + 1
+
+ if channel_last:
+ out_shape = (batch, out_height, out_width, out_channel)
+ else:
+ out_shape = (batch, out_channel, out_height, out_width)
+
+ return te.extern(
+ out_shape,
+ [src, weights],
+ lambda ins, outs: tvm.tir.call_packed(
+ "tvm.contrib.mkldnn.conv2d",
Review Comment:
I agree with that should be a unified name of DNNL, we already have USE_DNNL_CODEGEN and USE_MKLDNN for BYOC and `-libs`，I suggest to have USE_DNNL_LIBS for `-libs` and USE_DNNL_CODEGEN for BYOC, and change 'mkldnn' to 'dnnl' in codes.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] yangulei commented on a diff in pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
yangulei commented on code in PR #11571:
URL: https://github.com/apache/tvm/pull/11571#discussion_r890243914
##########
tests/python/relay/test_op_level2.py:
##########
@@ -1996,5 +1996,91 @@ def test_conv2d_rocm_sdot4():
np.testing.assert_equal(out, ref)
+@tvm.testing.requires_x86
+def test_conv2d_nchw_mkldnn():
+ d_shape = (1, 64, 56, 56)
+ w_shape = (64, 64, 3, 3)
+ padding = (1, 1)
+ strides = (1, 1)
+
+ data = relay.var("data", shape=d_shape, dtype="float32")
+ weight = relay.var("weight", shape=w_shape, dtype="float32")
+ out_channel = w_shape[0]
+ conv2d = relay.nn.conv2d(
+ data=data,
+ weight=weight,
+ kernel_size=w_shape[2:],
+ channels=out_channel,
+ padding=padding,
+ strides=strides,
+ out_dtype="float32",
+ )
+
+ mod = tvm.IRModule.from_expr(conv2d)
+
+ data_np = np.random.uniform(1, 10, d_shape).astype("float32")
+ weight_np = np.random.uniform(1, 10, size=w_shape).astype("float32")
+
+ target = "llvm -mcpu=skylake-avx512 -libs=mkldnn"
+ with tvm.transform.PassContext(opt_level=3):
+ lib = relay.build(mod, target=target, params={"weight": weight_np})
+
+ dev = tvm.device(target, 0)
+ runtime = tvm.contrib.graph_executor.GraphModule(lib["default"](dev))
+
+ runtime.set_input("data", data_np)
+ runtime.run()
+
+ out = runtime.get_output(0).numpy()
+
+ ref = tvm.topi.testing.conv2d_nchw_python(data_np, weight_np, strides, padding)
+
+ np.testing.assert_allclose(out, ref, rtol=1e-5, atol=1e-5)
+
+
+@tvm.testing.requires_x86
+def test_conv2d_nhwc_mkldnn():
+ d_shape = (1, 56, 56, 64)
+ w_shape = (3, 3, 64, 64)
+ padding = (1, 1)
+ strides = (1, 1)
+
+ data = relay.var("data", shape=d_shape, dtype="float32")
+ weight = relay.var("weight", shape=w_shape, dtype="float32")
+ out_channel = w_shape[3]
+ conv2d = relay.nn.conv2d(
+ data=data,
+ weight=weight,
+ kernel_size=w_shape[:2],
+ channels=out_channel,
+ padding=padding,
+ strides=strides,
+ out_dtype="float32",
+ data_layout="NHWC",
+ kernel_layout="HWIO",
+ )
+
+ mod = tvm.IRModule.from_expr(conv2d)
+
+ data_np = np.random.uniform(1, 10, d_shape).astype("float32")
+ weight_np = np.random.uniform(1, 10, size=w_shape).astype("float32")
+
+ target = "llvm -mcpu=skylake-avx512 -libs=mkldnn"
Review Comment:
Looks like this cause the failure in CI, since the CPU in CI doesn't support AVX512. We should check the ISA before using them as in [the tests for bfloat16](https://github.com/apache/tvm/pull/11111/commits/4caede47b08f6f5eded9e8d8bcc15da0edba4c68). It's not a robust or elegant way to do this, let me know if you have better idea.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] comaniac commented on a diff in pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
comaniac commented on code in PR #11571:
URL: https://github.com/apache/tvm/pull/11571#discussion_r891648302
##########
python/tvm/contrib/mkldnn.py:
##########
@@ -50,3 +51,107 @@ def matmul(lhs, rhs, transa=False, transb=False, **kwargs):
name="C",
**kwargs,
)
+
+
+def dnnl_conv2d(
+ src,
+ weights,
+ stride,
+ padding,
+ dilation,
+ groups,
+ channel_last=False,
+ out_dtype="float32",
+ **kwargs,
+):
+ """Convolution operator in NCHW layout.
+
+ Parameters
+ ----------
+ src : tvm.te.Tensor
+ 4-D with shape [batch, in_channel, in_height, in_width]
+
+ weights : tvm.te.Tensor
+ 4-D with shape [num_filter, in_channel, filter_height, filter_width]
+
+ stride : int or a list/tuple of two ints
+ Stride size, or [stride_height, stride_width]
+
+ padding : int or a list/tuple of 2 or 4 ints
+ padding size, or
+ [pad_height, pad_width] for 2 ints, or
+ [pad_top, pad_left, pad_bottom, pad_right] for 4 ints
+
+ dilation: int or a list/tuple of two ints
+ dilation size, or [dilation_height, dilation_width]
+
+ groups: str
+ input data layout: NCHW or NHWC
+
+ channel_last: bool
+ chose if input/output data format is in channel_last format(NHWC) or
+ in plain format(NCHW)
+
+ out_dtype: str
+ output datatype: now only support float32
+
+ Returns
+ -------
+ Output : tvm.te.Tensor
+ 4-D with shape [batch, out_channel, out_height, out_width]
+ """
+
+ assert isinstance(stride, int) or len(stride) == 2
+ assert isinstance(dilation, int) or len(dilation) == 2
+ if isinstance(stride, int):
+ stride_h = stride_w = stride
+ else:
+ stride_h, stride_w = stride
+
+ if isinstance(dilation, int):
+ dilation_h = dilation_w = dilation
+ else:
+ dilation_h, dilation_w = dilation
+
+ if channel_last:
+ batch, in_height, in_width, _ = src.shape
+ kernel_h, kernel_w, _, num_filter = weights.shape
+ else:
+ batch, _, in_height, in_width = src.shape
+ num_filter, _, kernel_h, kernel_w = weights.shape
+
+ dilated_kernel_h = (kernel_h - 1) * dilation_h + 1
+ dilated_kernel_w = (kernel_w - 1) * dilation_w + 1
+ pad_top, pad_left, pad_down, pad_right = get_pad_tuple(
+ padding, (dilated_kernel_h, dilated_kernel_w)
+ )
+ out_channel = num_filter
+ out_height = (in_height - dilated_kernel_h + pad_top + pad_down) // stride_h + 1
+ out_width = (in_width - dilated_kernel_w + pad_left + pad_right) // stride_w + 1
+
+ if channel_last:
+ out_shape = (batch, out_height, out_width, out_channel)
+ else:
+ out_shape = (batch, out_channel, out_height, out_width)
+
+ return te.extern(
+ out_shape,
+ [src, weights],
+ lambda ins, outs: tvm.tir.call_packed(
+ "tvm.contrib.mkldnn.conv2d",
Review Comment:
The name is a bit confusing...so does the use of libraries. We now have `USE_MKLDNN` (for cblas with matmul/dense) and `USE_DNNL` (for DNNL/OneDNN with matmal/dense/conv2d). AFAIK, MKL-DNN can be covered by DNNL, so should we deprecate MKL-DNN and use DNNL for both cases (e.g., `-libs` and BYOC)?
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] comaniac commented on pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
comaniac commented on PR #11571:
URL: https://github.com/apache/tvm/pull/11571#issuecomment-1146671027
Will take a look early next week.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] comaniac commented on a diff in pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
comaniac commented on code in PR #11571:
URL: https://github.com/apache/tvm/pull/11571#discussion_r891873046
##########
python/tvm/contrib/mkldnn.py:
##########
@@ -50,3 +51,107 @@ def matmul(lhs, rhs, transa=False, transb=False, **kwargs):
name="C",
**kwargs,
)
+
+
+def dnnl_conv2d(
+ src,
+ weights,
+ stride,
+ padding,
+ dilation,
+ groups,
+ channel_last=False,
+ out_dtype="float32",
+ **kwargs,
+):
+ """Convolution operator in NCHW layout.
+
+ Parameters
+ ----------
+ src : tvm.te.Tensor
+ 4-D with shape [batch, in_channel, in_height, in_width]
+
+ weights : tvm.te.Tensor
+ 4-D with shape [num_filter, in_channel, filter_height, filter_width]
+
+ stride : int or a list/tuple of two ints
+ Stride size, or [stride_height, stride_width]
+
+ padding : int or a list/tuple of 2 or 4 ints
+ padding size, or
+ [pad_height, pad_width] for 2 ints, or
+ [pad_top, pad_left, pad_bottom, pad_right] for 4 ints
+
+ dilation: int or a list/tuple of two ints
+ dilation size, or [dilation_height, dilation_width]
+
+ groups: str
+ input data layout: NCHW or NHWC
+
+ channel_last: bool
+ chose if input/output data format is in channel_last format(NHWC) or
+ in plain format(NCHW)
+
+ out_dtype: str
+ output datatype: now only support float32
+
+ Returns
+ -------
+ Output : tvm.te.Tensor
+ 4-D with shape [batch, out_channel, out_height, out_width]
+ """
+
+ assert isinstance(stride, int) or len(stride) == 2
+ assert isinstance(dilation, int) or len(dilation) == 2
+ if isinstance(stride, int):
+ stride_h = stride_w = stride
+ else:
+ stride_h, stride_w = stride
+
+ if isinstance(dilation, int):
+ dilation_h = dilation_w = dilation
+ else:
+ dilation_h, dilation_w = dilation
+
+ if channel_last:
+ batch, in_height, in_width, _ = src.shape
+ kernel_h, kernel_w, _, num_filter = weights.shape
+ else:
+ batch, _, in_height, in_width = src.shape
+ num_filter, _, kernel_h, kernel_w = weights.shape
+
+ dilated_kernel_h = (kernel_h - 1) * dilation_h + 1
+ dilated_kernel_w = (kernel_w - 1) * dilation_w + 1
+ pad_top, pad_left, pad_down, pad_right = get_pad_tuple(
+ padding, (dilated_kernel_h, dilated_kernel_w)
+ )
+ out_channel = num_filter
+ out_height = (in_height - dilated_kernel_h + pad_top + pad_down) // stride_h + 1
+ out_width = (in_width - dilated_kernel_w + pad_left + pad_right) // stride_w + 1
+
+ if channel_last:
+ out_shape = (batch, out_height, out_width, out_channel)
+ else:
+ out_shape = (batch, out_channel, out_height, out_width)
+
+ return te.extern(
+ out_shape,
+ [src, weights],
+ lambda ins, outs: tvm.tir.call_packed(
+ "tvm.contrib.mkldnn.conv2d",
Review Comment:
I don't think we really need to have 2 flags? It seems fine to enable both libs and BYOC when USE_DNNL is ON.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Qianshui-Jiang commented on pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
Qianshui-Jiang commented on PR #11571:
URL: https://github.com/apache/tvm/pull/11571#issuecomment-1148221868
@yangulei thanks for your suggesttion ! : ) and I've added the check for registed external function in test case, seems it works~
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] yangulei commented on a diff in pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
yangulei commented on code in PR #11571:
URL: https://github.com/apache/tvm/pull/11571#discussion_r890243914
##########
tests/python/relay/test_op_level2.py:
##########
@@ -1996,5 +1996,91 @@ def test_conv2d_rocm_sdot4():
np.testing.assert_equal(out, ref)
+@tvm.testing.requires_x86
+def test_conv2d_nchw_mkldnn():
+ d_shape = (1, 64, 56, 56)
+ w_shape = (64, 64, 3, 3)
+ padding = (1, 1)
+ strides = (1, 1)
+
+ data = relay.var("data", shape=d_shape, dtype="float32")
+ weight = relay.var("weight", shape=w_shape, dtype="float32")
+ out_channel = w_shape[0]
+ conv2d = relay.nn.conv2d(
+ data=data,
+ weight=weight,
+ kernel_size=w_shape[2:],
+ channels=out_channel,
+ padding=padding,
+ strides=strides,
+ out_dtype="float32",
+ )
+
+ mod = tvm.IRModule.from_expr(conv2d)
+
+ data_np = np.random.uniform(1, 10, d_shape).astype("float32")
+ weight_np = np.random.uniform(1, 10, size=w_shape).astype("float32")
+
+ target = "llvm -mcpu=skylake-avx512 -libs=mkldnn"
+ with tvm.transform.PassContext(opt_level=3):
+ lib = relay.build(mod, target=target, params={"weight": weight_np})
+
+ dev = tvm.device(target, 0)
+ runtime = tvm.contrib.graph_executor.GraphModule(lib["default"](dev))
+
+ runtime.set_input("data", data_np)
+ runtime.run()
+
+ out = runtime.get_output(0).numpy()
+
+ ref = tvm.topi.testing.conv2d_nchw_python(data_np, weight_np, strides, padding)
+
+ np.testing.assert_allclose(out, ref, rtol=1e-5, atol=1e-5)
+
+
+@tvm.testing.requires_x86
+def test_conv2d_nhwc_mkldnn():
+ d_shape = (1, 56, 56, 64)
+ w_shape = (3, 3, 64, 64)
+ padding = (1, 1)
+ strides = (1, 1)
+
+ data = relay.var("data", shape=d_shape, dtype="float32")
+ weight = relay.var("weight", shape=w_shape, dtype="float32")
+ out_channel = w_shape[3]
+ conv2d = relay.nn.conv2d(
+ data=data,
+ weight=weight,
+ kernel_size=w_shape[:2],
+ channels=out_channel,
+ padding=padding,
+ strides=strides,
+ out_dtype="float32",
+ data_layout="NHWC",
+ kernel_layout="HWIO",
+ )
+
+ mod = tvm.IRModule.from_expr(conv2d)
+
+ data_np = np.random.uniform(1, 10, d_shape).astype("float32")
+ weight_np = np.random.uniform(1, 10, size=w_shape).astype("float32")
+
+ target = "llvm -mcpu=skylake-avx512 -libs=mkldnn"
Review Comment:
Looks like this cause the failure in CI, since the CPU in CI doesn't support AVX512. We should check the ISA before using them as in [the tests for bfloat16](https://github.com/apache/tvm/pull/11111/commits/4caede47b08f6f5eded9e8d8bcc15da0edba4c68). It's not a robust or elegant way to do this, let me know if you have better idea.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Qianshui-Jiang commented on pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
Qianshui-Jiang commented on PR #11571:
URL: https://github.com/apache/tvm/pull/11571#issuecomment-1149981098
@comaniac hi，I change the cmake config flags and the symbol `mkldnn` to `dnnl`, now setup USE_DNNL can make both `-libs` and `dnnl_codegen` enabled.
BTW，seems a lot of files are influenced, shoud we splite this PR into 2 threads? pls take a look and comments more. thx a lot !
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi merged pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
masahi merged PR #11571:
URL: https://github.com/apache/tvm/pull/11571
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Qianshui-Jiang commented on pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
Qianshui-Jiang commented on PR #11571:
URL: https://github.com/apache/tvm/pull/11571#issuecomment-1146634479
@comaniac @AndrewZhaoLuo @junrushao1994 @masahi could you pls help review and have some suggestions on how to organized the file to USE_MKLDNN not only for cblas ? should the test case added in test_dnnl.py?
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Qianshui-Jiang commented on pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
Qianshui-Jiang commented on PR #11571:
URL: https://github.com/apache/tvm/pull/11571#issuecomment-1151706427
> The current scope seems good to me. You could either make this PR upstream compatible and rebase #11638, or rebase this PR after 11638 has merged. I'm fine with either way.
Seems #11638 was suggested to support `-libs=mkldnn` continuously but translated it to `-libs=dnnl`, litlle bit more tiny change would be considered,
Could you pls approvel merge this one at first? so that we could updating all of which related to `-libs=mkldnn` together in 11638.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Qianshui-Jiang commented on pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
Qianshui-Jiang commented on PR #11571:
URL: https://github.com/apache/tvm/pull/11571#issuecomment-1150756869
@comaniac I reset this PR to compatible with current mainline (use mkldnn), only contains the change described in title. or we can update them all in #11638 ?
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Qianshui-Jiang commented on pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
Qianshui-Jiang commented on PR #11571:
URL: https://github.com/apache/tvm/pull/11571#issuecomment-1150629148
@comaniac @masahi @areusch, hi guys，new PR created in #11638 , I would rebase this PR after that updated in mainline.
pls take a look, thx !
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] comaniac commented on pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
comaniac commented on PR #11571:
URL: https://github.com/apache/tvm/pull/11571#issuecomment-1150145654
It's a good idea. Please split the part that unifies MKL-DNN and DNNL to a separate PR.
Since it involves compatibility changes (e.g., `config.cmake`), it may need more reviews and visibility.
cc @masahi @areusch
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] comaniac commented on pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
comaniac commented on PR #11571:
URL: https://github.com/apache/tvm/pull/11571#issuecomment-1151400158
The current scope seems good to me. You could either make this PR upstream compatible and rebase #11638, or rebase this PR after 11638 has merged. I'm fine with either way.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Qianshui-Jiang commented on a diff in pull request #11571: [DNNL][Relay extern-schedule] DNNL Conv2D Kernel enable by assigning "-libs=mkldnn"
Posted by GitBox <gi...@apache.org>.
Qianshui-Jiang commented on code in PR #11571:
URL: https://github.com/apache/tvm/pull/11571#discussion_r891880282
##########
python/tvm/contrib/mkldnn.py:
##########
@@ -50,3 +51,107 @@ def matmul(lhs, rhs, transa=False, transb=False, **kwargs):
name="C",
**kwargs,
)
+
+
+def dnnl_conv2d(
+ src,
+ weights,
+ stride,
+ padding,
+ dilation,
+ groups,
+ channel_last=False,
+ out_dtype="float32",
+ **kwargs,
+):
+ """Convolution operator in NCHW layout.
+
+ Parameters
+ ----------
+ src : tvm.te.Tensor
+ 4-D with shape [batch, in_channel, in_height, in_width]
+
+ weights : tvm.te.Tensor
+ 4-D with shape [num_filter, in_channel, filter_height, filter_width]
+
+ stride : int or a list/tuple of two ints
+ Stride size, or [stride_height, stride_width]
+
+ padding : int or a list/tuple of 2 or 4 ints
+ padding size, or
+ [pad_height, pad_width] for 2 ints, or
+ [pad_top, pad_left, pad_bottom, pad_right] for 4 ints
+
+ dilation: int or a list/tuple of two ints
+ dilation size, or [dilation_height, dilation_width]
+
+ groups: str
+ input data layout: NCHW or NHWC
+
+ channel_last: bool
+ chose if input/output data format is in channel_last format(NHWC) or
+ in plain format(NCHW)
+
+ out_dtype: str
+ output datatype: now only support float32
+
+ Returns
+ -------
+ Output : tvm.te.Tensor
+ 4-D with shape [batch, out_channel, out_height, out_width]
+ """
+
+ assert isinstance(stride, int) or len(stride) == 2
+ assert isinstance(dilation, int) or len(dilation) == 2
+ if isinstance(stride, int):
+ stride_h = stride_w = stride
+ else:
+ stride_h, stride_w = stride
+
+ if isinstance(dilation, int):
+ dilation_h = dilation_w = dilation
+ else:
+ dilation_h, dilation_w = dilation
+
+ if channel_last:
+ batch, in_height, in_width, _ = src.shape
+ kernel_h, kernel_w, _, num_filter = weights.shape
+ else:
+ batch, _, in_height, in_width = src.shape
+ num_filter, _, kernel_h, kernel_w = weights.shape
+
+ dilated_kernel_h = (kernel_h - 1) * dilation_h + 1
+ dilated_kernel_w = (kernel_w - 1) * dilation_w + 1
+ pad_top, pad_left, pad_down, pad_right = get_pad_tuple(
+ padding, (dilated_kernel_h, dilated_kernel_w)
+ )
+ out_channel = num_filter
+ out_height = (in_height - dilated_kernel_h + pad_top + pad_down) // stride_h + 1
+ out_width = (in_width - dilated_kernel_w + pad_left + pad_right) // stride_w + 1
+
+ if channel_last:
+ out_shape = (batch, out_height, out_width, out_channel)
+ else:
+ out_shape = (batch, out_channel, out_height, out_width)
+
+ return te.extern(
+ out_shape,
+ [src, weights],
+ lambda ins, outs: tvm.tir.call_packed(
+ "tvm.contrib.mkldnn.conv2d",
Review Comment:
Yep, that would be more concise, 💯
I'll try and commit later.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org