Conv2d different interfaces are generated for different shape


#1
import numpy as np
import tvm
from tvm import autotvm
import topi
import topi.testing
from tvm.contrib.pickle_memoize import memoize
from topi.util import get_const_tuple
import argparse

parser = argparse.ArgumentParser()


def conv2d_nchw_codegen(batch, in_channel, in_height, in_width, num_filter, kernel_h, kernel_w, stride, padding,
                        function_name = "dsbdsfk", dilation=1):
    A = tvm.placeholder((batch, in_channel, in_height, in_width), name='input')
    W = tvm.placeholder((num_filter, in_channel, kernel_h, kernel_w), name='filter')

    a_shape = get_const_tuple(A.shape)
    w_shape = get_const_tuple(W.shape)
    dtype = A.dtype

    @memoize("topi.tests.test_topi_conv2d_nchw.verify_conv2d_nchw")
    def get_ref_data():
        a_np = np.random.uniform(size=a_shape).astype(dtype)
        w_np = np.random.uniform(size=w_shape).astype(dtype)
        dw_np = topi.testing.dilate_python(w_np, (1, 1, dilation, dilation))
        c_np = topi.testing.conv2d_nchw_python(a_np, dw_np, stride, padding)
        return a_np, w_np, c_np

    a_np, w_np, c_np = get_ref_data()

    # print(c_np)
    def check_device(device):
        ctx = tvm.context(device, 0)
        if not ctx.exist:
            print("Skip because %s is not enabled" % device)
            return
        with tvm.target.create(device):
            C = topi.nn.conv2d(A, W, (stride, stride), (padding, padding),
                               (dilation, dilation), layout='NCHW', out_dtype=dtype)
            s = topi.generic.schedule_conv2d_nchw([C])

        a = tvm.nd.array(a_np, ctx)
        w = tvm.nd.array(w_np, ctx)
        c = tvm.nd.array(np.zeros(get_const_tuple(C.shape), dtype=C.dtype), ctx)
        func = tvm.build(s, [A, W, C], device, name=function_name)
        func(a, w, c)
        dev_module = func.imported_modules[0]
        print(dev_module.get_source())
        tvm.testing.assert_allclose(c.asnumpy(), c_np, rtol=1e-4)

    for device in ["cuda"]:
        with autotvm.tophub.context(device):  # load tophub pre-tuned parameters
            check_device(device)

if __name__ == "__main__":
    #conv2d_nchw_codegen(1, 3, 224, 224, 64, 7, 7, 2, 3)
    conv2d_nchw_codegen(1, 64, 56, 56, 64, 3, 3, 1, 1)

when i call function: conv2d_nchw_codegen(1, 3, 224, 224, 64, 7, 7, 2, 3)
it will produce such a function:

extern "C" __global__ void dsbdsfk_kernel0( float* __restrict__ input,  float* __restrict__ filter,  float* __restrict__ compute)

I can use it with

dsbdsfk_kernel0(input, filter, out)

when i call function: conv2d_nchw_codegen(1, 64, 56, 56, 64, 3, 3, 1, 1)
t will produce such four functions:

extern "C" __global__ void dsbdsfk_kernel0( float* __restrict__ kernel_pack,  float* __restrict__ filter)
extern "C" __global__ void dsbdsfk_kernel1( float* __restrict__ input,  float* __restrict__ data_pack)
extern "C" __global__ void dsbdsfk_kernel2( float* __restrict__ kernel_pack,  float* __restrict__ data_pack,  float* __restrict__ bgemm)
extern "C" __global__ void dsbdsfk_kernel3( float* __restrict__ bgemm,  float* __restrict__ output)

No calls between functions.
I don’t know how to use the codegen.


#2

Maybe I can use
dsbdsfk_kernel0(kernel_pack, filter)
dsbdsfk_kernel1(input, data_pack)
dsbdsfk_kernel2(kernel_pack, data_pack, bgemm)
dsbdsfk_kernel3(bgemm, out)

But how size is the kernel_pack & data_pack & bgemm


#3

Can anyone help me? Thanks!