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.