def nhwc_tensorcore_cuda(cfg, Input, Filter, stride, padding, dilation, out_dtype):
    'Compute declaration for tensorcore'
    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
    (batch, in_height, in_width, in_channel) = get_const_tuple(Input.shape)
    (kernel_h, kernel_w, _, num_filter) = get_const_tuple(Filter.shape)
    assert ((((batch % 16) == 0) and ((in_channel % 16) == 0) and ((num_filter % 16) == 0)) or (((batch % 8) == 0) and ((in_channel % 16) == 0) and ((num_filter % 32) == 0)) or (((batch % 32) == 0) and ((in_channel % 16) == 0) and ((num_filter % 8) == 0))), 'The shape of (batch, in_channel, num_filter) must be multiple of (16, 16, 16) or (32, 16, 8) or (8, 16, 32) for now'
    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 = simplify((((((in_height - dilated_kernel_h) + pad_top) + pad_down) // stride_h) + 1))
    out_width = simplify((((((in_width - dilated_kernel_w) + pad_left) + pad_right) // stride_w) + 1))
    pad_before = [0, pad_top, pad_left, 0]
    pad_after = [0, pad_down, pad_right, 0]
    PaddedInput = pad(Input, pad_before, pad_after, name='PaddedInput')
    rc = te.reduce_axis((0, in_channel), name='rc')
    ry = te.reduce_axis((0, kernel_h), name='ry')
    rx = te.reduce_axis((0, kernel_w), name='rx')
    TransPaddedInput = te.compute(PaddedInput.shape, (lambda n, h, w, c: PaddedInput[(n, h, w, c)].astype('float16')))
    TransFilter = te.compute(Filter.shape, (lambda h, w, i, o: Filter[(h, w, i, o)].astype('float16')))
    Output = te.compute((batch, out_height, out_width, out_channel), (lambda nn, yy, xx, ff: te.sum((TransPaddedInput[(nn, ((yy * stride_h) + (ry * dilation_h)), ((xx * stride_w) + (rx * dilation_w)), rc)].astype(out_dtype) * TransFilter[(ry, rx, rc, ff)].astype(out_dtype)), axis=[ry, rx, rc])), name='Conv2dOutput', tag='conv2d_nhwc_tensorcore')
    return Output
Compute declaration for tensorcore
def nhwc_tensorcore_cuda(cfg, Input, Filter, stride, padding, dilation, out_dtype):
    
    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
    (batch, in_height, in_width, in_channel) = get_const_tuple(Input.shape)
    (kernel_h, kernel_w, _, num_filter) = get_const_tuple(Filter.shape)
    assert ((((batch % 16) == 0) and ((in_channel % 16) == 0) and ((num_filter % 16) == 0)) or (((batch % 8) == 0) and ((in_channel % 16) == 0) and ((num_filter % 32) == 0)) or (((batch % 32) == 0) and ((in_channel % 16) == 0) and ((num_filter % 8) == 0))), 'The shape of (batch, in_channel, num_filter) must be multiple of (16, 16, 16) or (32, 16, 8) or (8, 16, 32) for now'
    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 = simplify((((((in_height - dilated_kernel_h) + pad_top) + pad_down) // stride_h) + 1))
    out_width = simplify((((((in_width - dilated_kernel_w) + pad_left) + pad_right) // stride_w) + 1))
    pad_before = [0, pad_top, pad_left, 0]
    pad_after = [0, pad_down, pad_right, 0]
    PaddedInput = pad(Input, pad_before, pad_after, name='PaddedInput')
    rc = te.reduce_axis((0, in_channel), name='rc')
    ry = te.reduce_axis((0, kernel_h), name='ry')
    rx = te.reduce_axis((0, kernel_w), name='rx')
    TransPaddedInput = te.compute(PaddedInput.shape, (lambda n, h, w, c: PaddedInput[(n, h, w, c)].astype('float16')))
    TransFilter = te.compute(Filter.shape, (lambda h, w, i, o: Filter[(h, w, i, o)].astype('float16')))
    Output = te.compute((batch, out_height, out_width, out_channel), (lambda nn, yy, xx, ff: te.sum((TransPaddedInput[(nn, ((yy * stride_h) + (ry * dilation_h)), ((xx * stride_w) + (rx * dilation_w)), rc)].astype(out_dtype) * TransFilter[(ry, rx, rc, ff)].astype(out_dtype)), axis=[ry, rx, rc])), name='Conv2dOutput', tag='conv2d_nhwc_tensorcore')
    return Output