         outInfo, inInfo, reductionSize,                                 \
         (TYPE) outElements, init, modifyOp, reduceOp);                  \
   } else {                                                              \
-    kernelReduceNoncontigDim<ModifyOp, ReduceOp,                        \
-                             typename TensorUtils<TensorType>::DataType, \
-                             TYPE, OUT, IN>                             \
-      <<<grid, block, 0, THCState_getCurrentStream(state)>>>(           \
-        outInfo, inInfo, reductionStride, reductionSize,                \
+    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)                     \

  // 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,

    }

    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);
