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

void softmax_backward(int row, int col, const cnn::real *fx, const cnn::real *dEdf, cnn::real *dEdx)
{
    cudnnTensorDescriptor_t pInputDesc;
    int n = col; int c = 1; int h = 1; int w = row; 
    cnn::real one = 1.0;

    CHECK_CUDNN(cudnnCreateTensorDescriptor(&pInputDesc));
    CHECK_CUDNN(cudnnSetTensor4dDescriptor(pInputDesc, CUDNN_TENSOR_NCHW, cudnnDataType, n, c, h, w));
    CHECK_CUDNN(cudnnSoftmaxBackward(cudnn_handle, CUDNN_SOFTMAX_ACCURATE, CUDNN_SOFTMAX_MODE_INSTANCE,
        &one, pInputDesc, fx, pInputDesc, dEdf,

void vrelu_backward(int n, const cnn::real* fx, const cnn::real* dEdf, cnn::real* dEdx) {
  auto tb = SizeToBlockThreadPair(n);
  accBinaryExprKernel<<<tb.first, tb.second>>>(n, fx, dEdf, dEdx, FRectifyBackward());
}

void vexponential_linear_units(int n, const cnn::real* x, const cnn::real scale, cnn::real* y) {
    auto tb = SizeToBlockThreadPair(n);
    unaryExprKernel << <tb.first, tb.second >> >(n, x, y, FExponentialLinearUnits(scale));
}

void vexponential_linear_units_backward(int n, const cnn::real* fx, const cnn::real* dEdf, const cnn::real scale, cnn::real* dEdx) {
    auto tb = SizeToBlockThreadPair(n);
    accBinaryExprKernel << <tb.first, tb.second >> >(n, fx, dEdf, dEdx, FExponentialLinearUnitsBackward(scale));
}

void vtanh(int n, const cnn::real* x, cnn::real* y) {
