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

    auto tb = SizeToBlockThreadPair(col * row);
    accBinaryExprKernel << <tb.first, tb.second >> >(col * row, dEdf, gpu_softmax, dEdx, FSubtract());

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,
