From ecb6d10eec2386f0882a901be8cfedec071e93cb Mon Sep 17 00:00:00 2001 From: Thomas Viehmann Date: Mon, 30 Mar 2020 21:44:59 +0200 Subject: [PATCH 1/2] fix miopen convolutions --- tests/python/contrib/test_miopen.py | 8 +++----- topi/python/topi/rocm/conv2d.py | 4 +++- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/tests/python/contrib/test_miopen.py b/tests/python/contrib/test_miopen.py index b4bedd84e2e1..b5c214221f92 100644 --- a/tests/python/contrib/test_miopen.py +++ b/tests/python/contrib/test_miopen.py @@ -56,8 +56,7 @@ def test_conv2d(): yshape = [x.value for x in Y.shape] import topi - with tvm.target.create("rocm -libs=miopen"): - s = topi.generic.schedule_extern(Y) + s = te.create_schedule(Y.op) def verify(): ctx = tvm.rocm(0) @@ -68,9 +67,8 @@ def verify(): f(x, w, y) Y_ref = topi.nn.conv2d_nchw(X, W, (stride_h, stride_w), (pad_h, pad_w), (dilation_h, dilation_w)) - with tvm.target.rocm(): - s_ref = topi.generic.schedule_conv2d_nchw([Y_ref]) - f_ref = tvm.build(s_ref, [X, W, Y_ref], "rocm") + s_ref = te.create_schedule(Y_ref.op) + f_ref = tvm.build(s_ref, [X, W, Y_ref], "rocm", target_host="llvm") y_ref = tvm.nd.array(np.random.uniform(-1, 1, yshape).astype(np.float32), ctx) f_ref(x, w, y_ref) print("Max abs diff:", np.max(np.abs(y.asnumpy() - y_ref.asnumpy()))) diff --git a/topi/python/topi/rocm/conv2d.py b/topi/python/topi/rocm/conv2d.py index 713647e4ca8a..24500a661587 100644 --- a/topi/python/topi/rocm/conv2d.py +++ b/topi/python/topi/rocm/conv2d.py @@ -24,7 +24,7 @@ from ..nn.util import get_pad_tuple @autotvm.register_topi_compute("conv2d_nchw_miopen.rocm") -def conv2d_nchw_miopen(cfg, data, kernel, strides, padding, dilation, out_dtype='float32'): +def conv2d_nchw_miopen(cfg, data, kernel, strides, padding, dilation, layout='NCHW', out_dtype='float32'): """Conv2D operator for rocm backend. Parameters @@ -58,6 +58,8 @@ def conv2d_nchw_miopen(cfg, data, kernel, strides, padding, dilation, out_dtype= CO, CI, KH, KW = get_const_tuple(kernel.shape) N, _, H, W = get_const_tuple(data.shape) + assert layout == 'NCHW' + # handle dilation stride_h, stride_w = (strides, strides) if isinstance(strides, int) else strides pt, pl, pb, pr = get_pad_tuple(padding, (KH, KW)) From dc731f84fe862557c7a4313576603f6167671d14 Mon Sep 17 00:00:00 2001 From: Thomas Viehmann Date: Mon, 30 Mar 2020 22:23:26 +0200 Subject: [PATCH 2/2] fix overly long lines --- tests/python/contrib/test_miopen.py | 3 ++- topi/python/topi/rocm/conv2d.py | 3 ++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/tests/python/contrib/test_miopen.py b/tests/python/contrib/test_miopen.py index b5c214221f92..ed671e0c4810 100644 --- a/tests/python/contrib/test_miopen.py +++ b/tests/python/contrib/test_miopen.py @@ -66,7 +66,8 @@ def verify(): y = tvm.nd.array(np.random.uniform(-1, 1, yshape).astype(np.float32), ctx) f(x, w, y) - Y_ref = topi.nn.conv2d_nchw(X, W, (stride_h, stride_w), (pad_h, pad_w), (dilation_h, dilation_w)) + Y_ref = topi.nn.conv2d_nchw(X, W, (stride_h, stride_w), (pad_h, pad_w), + (dilation_h, dilation_w)) s_ref = te.create_schedule(Y_ref.op) f_ref = tvm.build(s_ref, [X, W, Y_ref], "rocm", target_host="llvm") y_ref = tvm.nd.array(np.random.uniform(-1, 1, yshape).astype(np.float32), ctx) diff --git a/topi/python/topi/rocm/conv2d.py b/topi/python/topi/rocm/conv2d.py index 24500a661587..4ee18775b938 100644 --- a/topi/python/topi/rocm/conv2d.py +++ b/topi/python/topi/rocm/conv2d.py @@ -24,7 +24,8 @@ from ..nn.util import get_pad_tuple @autotvm.register_topi_compute("conv2d_nchw_miopen.rocm") -def conv2d_nchw_miopen(cfg, data, kernel, strides, padding, dilation, layout='NCHW', out_dtype='float32'): +def conv2d_nchw_miopen(cfg, data, kernel, strides, padding, dilation, + layout='NCHW', out_dtype='float32'): """Conv2D operator for rocm backend. Parameters