Merge with HU Chi branch (Don't use this! It's an incomplete version)

1. Support X_INT dataType for Sum function.
2. Minor error fixed.

source/tensor/core/arithmetic/Sum.cpp

@@ -90,32 +90,10 @@ void _Sum(const XTensor * a, const XTensor * b, XTensor * c, DTYPE beta)
                 /* when c != a, OpenBLAS needs to copy a to c first. This operation
                  slow down the speed, so just use OpenBLAS when c == a */
 #if defined(USE_BLAS)
-                if( c == a){
+                if (c == a) {
                     AXPY(a->unitNum,beta,bp,1,cp,1);
-                } else{
-                     int num = a->unitNum;
-                        if (num % 4 == 0) {
-                            for (int i = 0; i < num; i += 4) {
-                                cp[i] = ap[i] + bp[i] * beta;
-                                cp[i + 1] = ap[i + 1] + bp[i + 1] * beta;
-                                cp[i + 2] = ap[i + 2] + bp[i + 2] * beta;
-                                cp[i + 3] = ap[i + 3] + bp[i + 3] * beta;
-                            }
-                        }
-                        else if (num % 2 == 0) {
-                            for (int i = 0; i < num; i += 2) {
-                                cp[i] = ap[i] + bp[i] * beta;
-                                cp[i + 1] = ap[i + 1] + bp[i + 1] * beta;
-                            }
-                        }
-                        else {
-                            for (int i = 0; i < num; i++) {
-                                cp[i] = ap[i] + bp[i] * beta;
-                            }
-                        }
                 }
-#else
-                    /* unrolling */
+                else {
                     int num = a->unitNum;
                     if (num % 4 == 0) {
                         for (int i = 0; i < num; i += 4) {
@@ -136,8 +114,61 @@ void _Sum(const XTensor * a, const XTensor * b, XTensor * c, DTYPE beta)
                             cp[i] = ap[i] + bp[i] * beta;
                         }
                     }
+                }
+#else
+                /* unrolling */
+                int num = a->unitNum;
+                if (num % 4 == 0) {
+                    for (int i = 0; i < num; i += 4) {
+                        cp[i] = ap[i] + bp[i] * beta;
+                        cp[i + 1] = ap[i + 1] + bp[i + 1] * beta;
+                        cp[i + 2] = ap[i + 2] + bp[i + 2] * beta;
+                        cp[i + 3] = ap[i + 3] + bp[i + 3] * beta;
+                    }
+                }
+                else if (num % 2 == 0) {
+                    for (int i = 0; i < num; i += 2) {
+                        cp[i] = ap[i] + bp[i] * beta;
+                        cp[i + 1] = ap[i + 1] + bp[i + 1] * beta;
+                    }
+                }
+                else {
+                    for (int i = 0; i < num; i++) {
+                        cp[i] = ap[i] + bp[i] * beta;
+                    }
+                }
 #endif
+            }
+            else if (a->dataType == X_INT &&
+                     b->dataType == X_INT &&
+                     c->dataType == X_INT)
+            {
+                int * ap = (int*)a->data;
+                int * bp = (int*)b->data;
+                int * cp = (int*)c->data;
+
+                /* unrolling */
+                int num = a->unitNum;
+                if (num % 4 == 0) {
+                    for (int i = 0; i < num; i += 4) {
+                        cp[i] = ap[i] + bp[i] * beta;
+                        cp[i + 1] = ap[i + 1] + bp[i + 1] * beta;
+                        cp[i + 2] = ap[i + 2] + bp[i + 2] * beta;
+                        cp[i + 3] = ap[i + 3] + bp[i + 3] * beta;
+                    }
                 }
+                else if (num % 2 == 0) {
+                    for (int i = 0; i < num; i += 2) {
+                        cp[i] = ap[i] + bp[i] * beta;
+                        cp[i + 1] = ap[i + 1] + bp[i + 1] * beta;
+                    }
+                }
+                else {
+                    for (int i = 0; i < num; i++) {
+                        cp[i] = ap[i] + bp[i] * beta;
+                    }
+                }
+            }
             else {
                 // TODO!!
                 ShowNTErrors("TODO!");

source/tensor/core/arithmetic/Sum.cu

@@ -45,6 +45,15 @@ void KernelADD(DTYPE * a, DTYPE * b, DTYPE * c, int size, DTYPE beta)
         c[i] = a[i] + b[i] * beta;
 }
 
+__global__
+void KernelADD(int * a, int * b, int * c, int size, int beta)
+{
+    int i = blockDim.x * blockIdx.x + threadIdx.x;
+
+    if (i < size)
+        c[i] = a[i] + b[i] * beta;
+}
+
 /*
 tensor summation c = a + b * \beta (cuda version)
 >> a - a tensor
@@ -100,6 +109,17 @@ void _CudaSum(const XTensor * a, const XTensor * b, XTensor * c, DTYPE beta)
                 KernelADD << <blocks, threads >> >((DTYPE*)a->data, (DTYPE*)b->data, (DTYPE*)c->data, a->unitNum, beta);
             }
         }
+        else if (a->dataType == X_INT &&
+                 b->dataType == X_INT &&
+                 c->dataType == X_INT)
+        {
+            int gridSize[3], blockSize[3];
+
+            GDevs.GetCudaThread(a->devID, a->unitNum, gridSize, blockSize);
+            dim3 blocks(gridSize[0]);
+            dim3 threads(blockSize[0]);
+            KernelADD << <blocks, threads >> >((int*)a->data, (int*)b->data, (int*)c->data, a->unitNum, (int)beta);
+        }
         else {
             // TODO!!
             ShowNTErrors("TODO!");