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

      sbuf[threadIdx.y][threadIdx.x] = init;
      unsigned col = bCol + threadIdx.x;
      if(row < size[1] && col < size[0]) {
        sbuf[threadIdx.y][threadIdx.x] = src[col];
      }
      __syncthreads();

      float* line = &sbuf[threadIdx.y][0];
      for(unsigned s = 16; s > 1; s >>= 1) {
        if(row < size[1] && threadIdx.x < s) {
          line[threadIdx.x] = op(line[threadIdx.x], line[threadIdx.x + s]);
        }
        __syncthreads();
      }
      if(reducing) {
        sbuf[threadIdx.x][0] = op(sbuf[threadIdx.x][0], sbuf[threadIdx.x][1]);
        acc = op(acc, sbuf[threadIdx.x][0]);
      }
      __syncthreads();
    }




template<class UnaryFunction, class BinaryFunction>
__host__ void THCudaTensor_transformReduceInnermostDim(THCudaTensor *tgt, THCudaTensor *src,
        UnaryFunction unary_op, BinaryFunction binary_op, float init)
{
  dim4 src_stride(0);
  dim4 tgt_stride(0);
