'Doing following operations (* is dot product)
C = alpha * A * B   + beta * C
C = alpha * A.T * B + beta * C
C = alpha * A * B.T + beta * C.
relu: if true applied before output (and prior to beta addition)
size: one of 32x128, 128x32, 64x128, 128x64, 128x128.  Sometimes the
fastest tiling isn\'t chosen for you.
Arguments:
A, B (GPUTensor): input operands
C (GPUTensor): output
alpha (float): scale A*B term
beta (float): scale C term before sum
relu (bool): whether to apply ReLu before output
size(nxm): Sometimes the fastest tiling isn\'t chosen for you.'

def compound_dot(self, A, B, C, alpha=1.0, beta=0.0, relu=False, bsum=None, repeat=1, size=None):

    assert (A.dtype.type == B.dtype.type == C.dtype.type)
    if (self.use_cudac_kernels or (B.shape[1] == 1)):
        for r in range(repeat):
            self.cublas_dot(A=A, B=B, C=C, alpha=alpha, beta=beta)
        if (bsum is not None):
            bsum[:] = self.sum(C, 1)
        return C
    assert (min(A.strides) == 1)
    assert (min(B.strides) == 1)
    assert (min(C.strides) == 1)
    lda = max(A.strides)
    ldb = max(B.strides)
    ldc = max(C.strides)
    if A.is_trans:
        opA = 't'
        if (size not in ('32x64', '16x64')):
            lda *= (8 * A.dtype.itemsize)
    else:
        opA = 'n'
    if B.is_trans:
        opB = 't'
    else:
        opB = 'n'
        if (size not in ('32x64', '16x64')):
            ldb *= (8 * B.dtype.itemsize)
    op = (opA + opB)
    assert (op != 'tt')
    m = A.shape[0]
    n = B.shape[1]
    k = A.shape[1]
    assert (m == C.shape[0])
    assert (n == C.shape[1])
    assert (k == B.shape[0])
    if (size is None):
        short = min(m, n)
        if (short < (384 - 16)):
            short128 = (short % 128)
            if (0 < short128 < 112):
                if (48 < short128 <= 64):
                    occupancy64 = (short // 64)
                    wide = max(m, n)
                    occupancy64 *= (((wide // 128) + ((wide % 128) != 0)) // _get_sm_count())
                    if (occupancy64 > 1):
                        size = 64
                    else:
                        size = 32
                else:
                    size = 32
            else:
                size = 128
        else:
            size = 128
        if (m >= n):
            if (op == 'nt'):
                size = 128
            (sizeA, sizeB) = (128, size)
        else:
            if (op == 'tn'):
                size = 128
            elif (size == 64):
                size = 32
            (sizeA, sizeB) = (size, 128)
        size = ('%dx%d' % (sizeA, sizeB))
    else:
        (sizeA, sizeB) = (int(s) for s in size.split('x'))
    gridA = ((m // sizeA) + ((m % sizeA) != 0))
    gridB = ((n // sizeB) + ((n % sizeB) != 0))
    k_vec = (8 if ((sizeA in (16, 32)) or (sizeB == 32)) else 16)
    vec_opt = None
    if (op == 'tn'):
        if (((m % 4) == 0) and ((n % 4) == 0) and ((A.strides[1] % 4) == 0) and ((B.strides[0] % 4) == 0)):
            vec_opt = ('vec',)
    elif (op == 'nn'):
        if (((k % k_vec) == 0) and ((n % 4) == 0) and ((A.strides[0] % k_vec) == 0) and ((B.strides[0] % 4) == 0)):
            vec_opt = ('vec',)
    elif (op == 'nt'):
        if (((k % k_vec) == 0) and ((n % 4) == 0) and ((A.strides[0] % k_vec) == 0) and ((B.strides[1] % k_vec) == 0)):
            vec_opt = ('vec',)
    if (C.dtype.type is np.float16):
        clss = 'hgemm'
    elif (C.dtype.type is np.float32):
        clss = 'sgemm'
    else:
        raise TypeError('Only    floating    point    dot    currently    supported.')
    flags = 0
    if relu:
        flags |= 2
    kernel = kernel_specs.get_kernel('_'.join((clss, op, size)), vec_opt)
    params = [(1, int(gridA), int(gridB)), (kernel.threads, 1, 1), self.stream, C.gpudata, A.gpudata, B.gpudata, alpha, beta, flags, int(lda), int(ldb), int(ldc), int(m), int(n), int(k), 0, 0, 0, 0]
    if (repeat > 1):
        for r in range(max((repeat // 10), 1)):
            kernel.prepared_async_call(*params)
    if (self.bench or (repeat > 1)):
        (start, end) = _get_events()
        start.record(self.stream)
    for r in range(repeat):
        kernel.prepared_async_call(*params)
    if (self.bench or (repeat > 1)):
        end.record(self.stream)
        end.synchronize()
        msecs = (end.time_since(start) / repeat)
        gflops = ((((m * n) * k) * 2.0) / (msecs * 1000000.0))
        neon_logger.display(('%7.3f    msecs    %4.0f    gflops    (%s_%s:    %d,%d,%d)    size:%s    grid:(%d,%d)' % (msecs, gflops, clss, op, m, n, k, size, gridA, gridB)))
        if (repeat > 1):
            return (msecs, gflops)
    if (bsum is not None):
        bsum[:] = self.sum(C, 1)
    return C
