  *
  * Reduction along other dimensions is handled in a separate kernel.
  */
-template<class Op>
-__global__ void THCudaTensor_kernel_reduceInnermostDim(float *tgt, float *src_,
-        dim4 src_stride, dim4 tgt_stride, dim4 size, Op op, float init)
+template<class UnaryFunction, class BinaryFunction>
+__global__ void THCudaTensor_kernel_transformReduceInnermostDim(float *tgt, float *src_,
+        dim4 src_stride, dim4 tgt_stride, dim4 size, UnaryFunction unary_op, BinaryFunction binary_op, float init)
 {
   __shared__ float sbuf[16][32]; // 8kB
 

    float *src = src_ + z * src_stride[2] + y * src_stride[1] + col;
    float acc = init;
    for(unsigned i=0; i < size[reduce]; i++) {
      acc = op(acc, *src);
      src += src_stride[reduce];
    }
    tgt[z * tgt_stride[2] + y * tgt_stride[1] + col] = float(acc);




template<class UnaryFunction, class BinaryFunction>
__host__ void THCudaTensor_transformReduceOuterDim(THCudaTensor *tgt, THCudaTensor *src,
        long rdim, UnaryFunction unary_op, BinaryFunction binary_op, float init)
{
  const size_t reduce = 3;
  dim4 src_stride(0);
