                                      BinaryFunction binary_op)
 {
   THArgCheck(dimension >= 0 && dimension < THCudaTensor_nDimension(state, src), 3, "dimension out of range");
-  THArgCheck(THCudaTensor_nDimension(state, src) <= 4, 2, "too many dimensions (>4)");
 
   THLongStorage *dim = THCudaTensor_newSizeOf(state, src);
   THLongStorage_set(dim, dimension, 1);

                                                   long rdim, thrust::pair<float,float> init,
                                                   BinaryFunction binary_op)
{
  const size_t reduce = 3;
  dim4 src_stride(0);
  dim4 tgt1_stride(0);
  dim4 tgt2_stride(0);
  dim4 size(1);

  unsigned ndim = THCudaTensor_nDimension(state, src);
  for(unsigned idim=0, o=ndim-2; idim < ndim; idim++) {
    unsigned odim = idim == rdim ? reduce : o--;
    src_stride[odim] = THCudaTensor_stride(state, src, idim);
    tgt1_stride[odim] = THCudaTensor_stride(state, tgt1, idim);
    tgt2_stride[odim] = THCudaTensor_stride(state, tgt2, idim);
    size[odim]       = THCudaTensor_size(state, src, idim);
  }

  const unsigned nThreadPerBlock = 256;
  unsigned nBlockPerColumn = DIVUP(size[0], nThreadPerBlock);
  dim3 threads(nThreadPerBlock);
  unsigned maxGridDim = 1024; // anything < 64k is fine. The choice has no impact on performance.
  dim3 grid(min(maxGridDim, nBlockPerColumn), min(maxGridDim, size[1]), min(maxGridDim, size[2]));

  THCudaTensor_kernel_transformReduceOuterDimIndex<<<grid, threads>>>(
    THCudaTensor_data(state, tgt1), THCudaTensor_data(state, tgt2),
    THCudaTensor_data(state, src), src_stride, tgt1_stride, tgt2_stride, size, init, binary_op);
  cudaError errcode = cudaGetLastError();
  if(errcode != cudaSuccess) {
    THError(cudaGetErrorString(errcode));

        long dimension, UnaryFunction unary_op, float init, BinaryFunction binary_op)
{
  THArgCheck(dimension >= 0 && dimension < THCudaTensor_nDimension(state, src), 3, "dimension out of range");

  THLongStorage *dim = THCudaTensor_newSizeOf(state, src);
  THLongStorage_set(dim, dimension, 1);
