     __global__ void _rowElementMultiplyWith(
         ElemType* us,
         const ElemType* a,
-        const int N, // us.GetNumRows();
-        const int M) // a.GetNumCols();
+        const int N, // usrow;
+        const int M) // acol;
     {
         int id = blockDim.x * blockIdx.x + threadIdx.x;
         if (id >= M)
             return;
 
-        // __shared__ ElemType _a[GridDim::maxThreadsPerBlock];
+        // __shared__ ElemType _a[MAX_THREADS_PER_BLOCK];
         // _a[threadIdx.x]=a[id];
         ElemType mul = a[id];
         for (int i = 0; i < N; ++i)

#pragma once

#include "cnn/macros.h"
namespace cnn {
namespace gpu {


        }
    }

    template <class ElemType>
    void AssignElementProductOf(int arow, int acol, const ElemType* a,
        int brow, int bcol, const ElemType* b,
        int crow, int ccol, ElemType* c)
    {
        assert(arow == brow && acol == bcol);
        if (!(arow == brow && acol == bcol))
            throw std::invalid_argument("The input matrix dimensions do not match.");

        if (!(crow == brow && ccol == bcol))
            throw std::invalid_argument("The input matrix dimensions do not match.");

        int N = (int)crow * ccol;
        int blocksPerGrid = (int)ceil(((double)N) / MAX_THREADS_PER_BLOCK);

        cudaEvent_t done = nullptr;
        cudaEventCreate(&done);
        _assignElementProductOf<ElemType> << <blocksPerGrid, MAX_THREADS_PER_BLOCK, 0, cudaStreamDefault >> >(c, a, b, N);
        cudaEventRecord(done);
        cudaEventSynchronize(done);
        cudaEventDestroy(done);
    }

    template <class ElemType>
    void InnerProduct(unsigned arow, unsigned acol, const ElemType* a, 
        unsigned brow, unsigned bcol, const ElemType* b, 
        unsigned crow, unsigned ccol, ElemType* c, 
        const bool isColWise)
    {
        const int m = (int)arow; 
        const int n = (int)acol;
        const int k = (int)brow;
        const int l = (int)bcol;

        assert(m > 0 && n > 0 && k > 0 && l > 0); // converting from size_t to int may cause overflow
        assert(m == k && n == l);                 // converting from size_t to int may cause overflow
        if (m != k || n != l)
            throw std::invalid_argument("Matrices a and b should have same dimension.");

        if (isColWise)
        {
            if (crow != 1 || ccol != n)
                throw std::invalid_argument("InnerProduct output matrix dimension mismatch");
        }
        else
        {
            if (crow != m || ccol != 1)
                throw std::invalid_argument("InnerProduct output matrix dimension mismatch");
        }

        if ((isColWise && m == 1) || !isColWise && n == 1) // in this case it's equivalent to element-wise product
        {
            AssignElementProductOf<ElemType>(arow, acol, a, brow, bcol, b, crow, ccol, c);
        }
        else
        {
            cudaEvent_t done = nullptr;

            int blocksPerGrid = 0;
            if (isColWise) // col-wise
            {
                blocksPerGrid = (int)ceil(1.0 * n / MAX_THREADS_PER_BLOCK);
            }
            else
            {
                blocksPerGrid = (int)ceil(1.0 * m / MAX_THREADS_PER_BLOCK);
            }

            cudaEventCreate(&done);
            _innerProduct<ElemType> << <blocksPerGrid, MAX_THREADS_PER_BLOCK, 0, cudaStreamDefault >> >(c, a, b, m, n, isColWise);
            cudaEventRecord(done);
            cudaEventSynchronize(done);
            cudaEventDestroy(done);
        }
    }

    /// <summary>Matrix-scalar multiply with col-major matrices: c = alpha * a + b</summary>
    /// if a is a column vector, add to all columns of b
    /// if a is a row vector, add to all rows of b
    /// if a is a scalar, add to all elements of b
    /// <param name="alpha">Scalar</param>
    /// <param name="a">Input matrix</param>
    /// <param name="b">Input matrix</param>
    /// <param name="c">Resulting matrix, user is responsible for allocating this</param>
    template <class ElemType>
    void ScaleAndAdd(ElemType alpha, int arow, int acol, const ElemType* a, int brow, int bcol, const ElemType* b,
        int crow, int ccol, ElemType* c)
    {
        return;

        if (crow != brow || ccol != bcol)
            throw std::invalid_argument("output dimension mismatch");

        if (arow == brow && acol == bcol) // dimensions match
        {

            unsigned N = crow * ccol;
            int blocksPerGrid = (int)ceil(1.0 * N / MAX_THREADS_PER_BLOCK);

            cudaEvent_t done = nullptr;
            cudaEventCreate(&done);
            _matrixMatrixAddOnCuda<ElemType> << <blocksPerGrid, MAX_THREADS_PER_BLOCK, 0, cudaStreamDefault >> >(alpha, a, b, c, N);
            cudaEventRecord(done);
            cudaEventSynchronize(done);
            cudaEventDestroy(done);
        }
        else if (arow * acol == 1)
        {
            unsigned N = crow * ccol;
            int blocksPerGrid = (int)ceil(1.0 * N / MAX_THREADS_PER_BLOCK);

            cudaEvent_t done = nullptr;
            cudaEventCreate(&done);
            _scaleAndAddScalar<ElemType> << <blocksPerGrid, MAX_THREADS_PER_BLOCK, 0, cudaStreamDefault >> >(c, N, alpha, a, b);
            cudaEventRecord(done);
            cudaEventSynchronize(done);
            cudaEventDestroy(done);
        }
        else if (acol == 1) // col vector, add it to all columns
        {
            unsigned m = (int)crow;
            unsigned n = (int)ccol;
            if (m != (int)arow)
                throw std::invalid_argument("To add column vector, rows should match.");

            cudaEvent_t done = nullptr;
            int blocksPerGrid = (int)(ceil(1.0 * m * n / MAX_THREADS_PER_BLOCK));
            cudaEventCreate(&done);
            _matrixVectorColumnWiseAddWithThreadPerElem<ElemType> << <blocksPerGrid, MAX_THREADS_PER_BLOCK, 0, cudaStreamDefault >> >(a, b, c, alpha, m, n);

            cudaEventRecord(done);
            cudaEventSynchronize(done);
            cudaEventDestroy(done);
        }
        else if (arow == 1) // row vector, add it to all rows
        {
            int m = (int)crow;
            int n = (int)ccol;
            if (n != (int)acol)
                throw std::invalid_argument("To add row vector, columns should match.");
            cudaEvent_t done = nullptr;
            int blocksPerGrid = (int)(ceil(1.0 * m * n / MAX_THREADS_PER_BLOCK));
            cudaEventCreate(&done);
            _matrixVectorRowWiseAddWithThreadPerElem<ElemType> << <blocksPerGrid, MAX_THREADS_PER_BLOCK, 0, cudaStreamDefault >> >(a, b, c, alpha, m, n);

            cudaEventRecord(done);
            cudaEventSynchronize(done);
            cudaEventDestroy(done);
        }
        else
            throw std::invalid_argument("dimension of matrix c does not match dimension of matrix a.");
    }

} // namespace gpu
} // namespace cnn

