fix reduceMax and reduceSum bug

source/tensor/core/reduce/ReduceMax.cu

@@ -544,11 +544,13 @@ void _CudaReduceMax(const XTensor * input, XTensor * output, int dim)
         dim3 grids;
         dim3 grids;
         dim3 blocks;
         dim3 blocks;
         continuousStorageThreadAllocation(grids, blocks, (long long)blockNum, strideNum);
         continuousStorageThreadAllocation(grids, blocks, (long long)blockNum, strideNum);
-        if (blocks.y > 128) {
+        if (blocks.y >= 128) {
             KernelReduceMaxOp <<<grids, blocks >>> ((DTYPE *)input->data, (DTYPE*)output->data, stride, strideNum, grids.y, blockSize, blockNum);
             KernelReduceMaxOp <<<grids, blocks >>> ((DTYPE *)input->data, (DTYPE*)output->data, stride, strideNum, grids.y, blockSize, blockNum);
         }
         }
         else {
         else {
-            KernelReduceMaxOpLessBlocks <<<blockNum / 4, 128 >>> ((DTYPE *)input->data, (DTYPE*)output->data, strideNum, blockNum);
+            if (blockNum % 4 != 0) blockNum = (int)(blockNum / 4) + 1;
+            else blockNum = blockNum / 4;
+            KernelReduceMaxOpLessBlocks <<<blockNum, 128 >>> ((DTYPE *)input->data, (DTYPE*)output->data, strideNum, blockNum);
         }
         }
     }
     }
     else {
     else {

source/tensor/core/reduce/ReduceSum.cu

@@ -730,10 +730,13 @@ void _CudaReduceSum(const XTensor * input, XTensor * output, int dim, const XTen
         dim3 grids;
         dim3 grids;
         dim3 blocks;
         dim3 blocks;
         continuousStorageThreadAllocation(grids, blocks, (long long)blockNum, strideNum);
         continuousStorageThreadAllocation(grids, blocks, (long long)blockNum, strideNum);
-        if (blocks.y > 128)
+        if (blocks.y >= 128)
             KernelReduceSumOp <<<grids, blocks >>> ((DTYPE *)input->data, (DTYPE*)output->data, stride, strideNum, grids.y, blockSize, blockNum, sp, power, isExp);
             KernelReduceSumOp <<<grids, blocks >>> ((DTYPE *)input->data, (DTYPE*)output->data, stride, strideNum, grids.y, blockSize, blockNum, sp, power, isExp);
-        else
-            KernelReduceSumOpLessBlocks <<<blockNum / 4, 128 >>> ((DTYPE *)input->data, (DTYPE*)output->data, strideNum, blockNum, sp, power, isExp);
+        else {
+            if (blockNum % 4 != 0) blockNum = (int)(blockNum / 4) + 1;
+            else blockNum = blockNum / 4;
+            KernelReduceSumOpLessBlocks << <blockNum, 128 >> > ((DTYPE *)input->data, (DTYPE*)output->data, strideNum, blockNum, sp, power, isExp);
+        }
     }
     }
     else if (stride != 1 && stride * blockNum > 4096){
     else if (stride != 1 && stride * blockNum > 4096){
         //GDevs->GetGridAndBlockSize2D(devID, stride * blockNum, strideNum,MAX_INT, cudaGridSize, cudaBlockSize);
         //GDevs->GetGridAndBlockSize2D(devID, stride * blockNum, strideNum,MAX_INT, cudaGridSize, cudaBlockSize);