     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);
+      acc = binary_op(acc, unary_op(*src));
       src += src_stride[reduce];
     }
     tgt[z * tgt_stride[2] + y * tgt_stride[1] + col] = float(acc);

 *
 * Reduction along the innermost dimension is handled in a separate kernel.
 */
template<class Op>
__global__ void THCudaTensor_kernel_reduceOuterDim(float *tgt, float *src_,
        dim4 src_stride, dim4 tgt_stride, dim4 size, Op op, float init)
{
  const size_t reduce = 3;


 *
 * Reduction along other dimensions is handled in a separate kernel.
 */
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

