(self, A, B, C, alpha=1.0, beta=0.0, relu=False, bsum=None,
                     repeat=1, size=None)
return C
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.
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.
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 .
def compound_dot(self, A, B, C, alpha=1.0, beta=0.0, relu=False, bsum=None,
                     repeat=1, size=None):

        """
        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.
        """
        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

        # one dimention must be contiguous
        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  # saves a kernel register
        else:
            opA = 'n'

        if B.is_trans:
            opB = 't'
        else:
            opB = 'n'
            if size not in ("32x64", "16x64"):
                ldb *= 8 * B.dtype.itemsize  # saves a kernel register

        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]

        # Some basic tile size selection.
        # Your best bet is to benchmark your code with all 3 sizes
        # and manually fine tune the selection for each layer.
        # TODO: Perhaps I'll add an autotuning mode.
        if size is None:
            # find the shorter side
            short = min(m, n)
            # anything bigger than this just use 128
            if short < 384 - 16:
                # compute remainder of 128
                short128 = short % 128
                # if remainder is more than 112 just use 128
                if 0 < short128 < 112:
                    # to figure out when to use 64 over 32 we need to calc
                    # occupancy at 64
                    if 48 < short128 <= 64:
                        occupancy64 = short // 64
                        wide = max(m, n)
                        occupancy64 *= (wide // 128 + (wide %
                                                       128 != 0)) // _get_sm_count()
                        # 64 is only faster than 32 when occupancy is more than
                        # 1 warp per scheduler.
                        if occupancy64 > 1:
                            size = 64
                        else:
                            size = 32
                    else:
                        size = 32
                else:
                    size = 128
            # There's a large regime where 64 is faster, but it's hard to
            # characterize
            else:
                size = 128

            # match the kernel to the optimal short size but avoid not
            # implemented kernels
            if m >= n:
                if op == "nt":
                    size = 128
                sizeA, sizeB = (128, size)
            else:
                if op == "tn":
                    size = 128
                # temp till I can write these kernels (coming soon)
                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",)

        # nt and nn are more efficient with k%16==0
        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]

        # Warmup
        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
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 # one dimention must be contiguous 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 # saves a kernel register else : opA = 'n' if B . is_trans : opB = 't' else : opB = 'n' if size not in ( "32x64" , "16x64" ) : ldb *= 8 * B . dtype . itemsize # saves a kernel register 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 ] # Some basic tile size selection. # Your best bet is to benchmark your code with all 3 sizes # and manually fine tune the selection for each layer. # TODO: Perhaps I'll add an autotuning mode. if size is None : # find the shorter side short = min ( m , n ) # anything bigger than this just use 128 if short < 384 - 16 : # compute remainder of 128 short128 = short % 128 # if remainder is more than 112 just use 128 if 0 < short128 < 112 : # to figure out when to use 64 over 32 we need to calc # occupancy at 64 if 48 < short128 <= 64 : occupancy64 = short // 64 wide = max ( m , n ) occupancy64 *= ( wide // 128 + ( wide % 128 != 0 ) ) // _get_sm_count ( ) # 64 is only faster than 32 when occupancy is more than # 1 warp per scheduler. if occupancy64 > 1 : size = 64 else : size = 32 else : size = 32 else : size = 128 # There's a large regime where 64 is faster, but it's hard to # characterize else : size = 128 # match the kernel to the optimal short size but avoid not # implemented kernels if m >= n : if op == "nt" : size = 128 sizeA , sizeB = ( 128 , size ) else : if op == "tn" : size = 128 # temp till I can write these kernels (coming soon) 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" , ) # nt and nn are more efficient with k%16==0 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 ] # Warmup 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