 void logsoftmax_backward(int row, int col, const cnn::real *fx, const cnn::real *dEdf, cnn::real *dEdx, cnn::real * gpu_softmax, cnn::real *grd)
 {
     vexp(row * col, fx, gpu_softmax);
-    VectorSum(row, col, dEdf, grd, true); 
-    RowElementMultiplyWith(1, col, grd, row, col, gpu_softmax);
+    vector_sum(row, col, dEdf, grd, true);
+    row_element_multiply_with(1, col, grd, row, col, gpu_softmax);
 
     auto tb = SizeToBlockThreadPair(col * row);
     accBinaryExprKernel << <tb.first, tb.second >> >(col * row, dEdf, gpu_softmax, dEdx, FSubtract());

void softmax(int row, int col, const cnn::real* x0, cnn::real* y)
{
    cudnnTensorDescriptor_t pInputDesc;
    int n = col; int c = 1; int h = 1; int w = row;

    cnn::real one = 1.0, zero = 0.0;
    CHECK_CUDNN(cudnnCreateTensorDescriptor(&pInputDesc));
    CHECK_CUDNN(cudnnSetTensor4dDescriptor(pInputDesc, CUDNN_TENSOR_NCHW, cudnnDataType, n, c, h, w));
    CHECK_CUDNN(cudnnSoftmaxForward(cudnn_handle, CUDNN_SOFTMAX_ACCURATE, CUDNN_SOFTMAX_MODE_INSTANCE,
        &one, pInputDesc, x0,

    unaryExprKernel << <tb.first, tb.second >> >(n, x, y, FConstantMinus(c));
}

void vconstant_minusx_backward(int n, cnn::real c, const cnn::real* x, cnn::real* y) {
    auto tb = SizeToBlockThreadPair(n);
    accUnaryExprKernel << <tb.first, tb.second >> >(n, x, y, FConstantMinus(c));
}

void vconstant_multiplyx(int n, cnn::real c, const cnn::real* x, cnn::real* y) {
    auto tb = SizeToBlockThreadPair(n);
    unaryExprKernel << <tb.first, tb.second >> >(n, x, y, FConstantMultiply(c));
