fix miopen pad

Author: t-vi

python/tvm/relay/op/strategy/rocm.py

@@ -36,6 +36,7 @@ def conv2d_strategy_rocm(attrs, inputs, out_type, target):
     layout = attrs.data_layout
     stride_h, stride_w = attrs.get_int_tuple("strides")
     kernel_layout = attrs.kernel_layout
+    padding = attrs.get_int_tuple("padding")
     if dilation_h < 1 or dilation_w < 1:
         raise ValueError("dilation should be positive value")
 
@@ -77,7 +78,8 @@ def conv2d_strategy_rocm(attrs, inputs, out_type, target):
         else:
             raise RuntimeError("Unsupported conv2d layout {} for CUDA".format(layout))
         # add miopen implementation
-        if "miopen" in target.libs and layout == "NCHW":
+        if "miopen" in target.libs and layout == "NCHW" and padding[0] == padding[2] and \
+            padding[1] == padding[3]:
             strategy.add_implementation(
                 wrap_compute_conv2d(topi.rocm.conv2d_nchw_miopen, True),
                 wrap_topi_schedule(topi.rocm.schedule_conv2d_nchw_miopen),

src/target/llvm/codegen_amdgpu.cc

@@ -219,8 +219,8 @@ runtime::Module BuildAMDGPU(IRModule mod, std::string target) {
          << " -mattr=-code-object-v3 "
          << target.substr(4, target.length() - 4);
   std::unique_ptr<llvm::TargetMachine> tm = GetLLVMTargetMachine(config.str());
-  std::unique_ptr<CodeGenAMDGPU> cg(new CodeGenAMDGPU());
   std::unique_ptr<llvm::LLVMContext> ctx(new llvm::LLVMContext());
+  std::unique_ptr<CodeGenAMDGPU> cg(new CodeGenAMDGPU());
 
   cg->Init("TVMAMDGPUModule", tm.get(), ctx.get(), false, false);
 

topi/python/topi/rocm/conv2d.py

@@ -66,7 +66,7 @@ def conv2d_nchw_miopen(cfg, data, kernel, strides, padding, dilation,
     pt, pl, pb, pr = get_pad_tuple(padding, (KH, KW))
     pad_h, pad_w = pt + pb, pl + pr
     dilation_h, dilation_w = (dilation, dilation) if isinstance(dilation, int) else dilation
-
+    assert (pt == pb) and (pl == pr)
     OH = (H + 2 * pad_h - KH) // stride_h + 1
     OW = (W + 2 * pad_w - KW) // stride_w + 1
     cfg.add_flop(2 * N * OH * OW * CO * CI * ((KH - 1) * dilation_h + 1) *\
@@ -76,8 +76,8 @@ def conv2d_nchw_miopen(cfg, data, kernel, strides, padding, dilation,
                                  kernel,
                                  stride_h,
                                  stride_w,
-                                 pad_h,
-                                 pad_w,
+                                 pt,
+                                 pl,
                                  dilation_h,
                                  dilation_w,
                                  conv_mode=0,