     }
 
     block = getNoncontigReduceBlock();
-  }
-
+
+    if(outElements <= 4096){
+        //x dim does different columns
+        //y dim helps with the same reduction
+        //If we only have 8 loops, don't bother sharing work across ydim
+        unsigned long ydim = THCCeilDiv(reductionSize, 8L);
+
+        //don't want y dim any bigger than 16, leaving min x dim to 32
+        ydim = min((unsigned long) 16, ydim);
+
+        block = dim3(THC_NONCONTIG_REDUCE_BLOCK_SIZE, 1, 1);
+        while(ydim > 1){
+          block.x /= 2;
+          block.y *= 2;
+          ydim /= 2;
+        }
+        THC_getGridFromTiles(THCCeilDiv(outElements, (long)block.x), grid);
+
+    }
+  }
   // Resize out to correspond to the reduced size
   THLongStorage* sizes = TensorUtils<TensorType>::newSizeOf(state, in);
   THLongStorage_set(sizes, dim, 1);

  // Each thread handles one slice
  return getLinearBlockId<IndexType>() * THC_NONCONTIG_REDUCE_BLOCK_SIZE + threadIdx.x;
}

// Kernel that handles an entire reduction of a slice of a tensor per each thread
template <typename ModifyOp,

        outInfo, inInfo, reductionSize,                                 \
        (TYPE) outElements, init, modifyOp, reduceOp);                  \
  } else {                                                              \
    if(block.y == 1){                                                   \
        kernelReduceNoncontigDim<ModifyOp, ReduceOp,                    \
                           typename TensorUtils<TensorType>::DataType,  \
                           TYPE, OUT, IN>                               \
        <<<grid, block, 0, THCState_getCurrentStream(state)>>>(         \
                       outInfo, inInfo, reductionStride, reductionSize, \
        (TYPE) outElements, init, modifyOp, reduceOp);                  \
    }else{                                                              \
        kernelReduceNoncontigDim_shared<ModifyOp, ReduceOp,             \
                           typename TensorUtils<TensorType>::DataType,  \
                           TYPE, OUT, IN>                               \
        <<<grid, block, 0, THCState_getCurrentStream(state)>>>(         \
                       outInfo, inInfo, reductionStride, reductionSize, \
                       (TYPE) outElements, init, modifyOp, reduceOp);   \
    }                                                                   \
  }                                                                     \

#define HANDLE_IN_CASE(TYPE, OUT, IN)                     \
