Support generating code for grad_op

[Xreki]

• 增加Operation和OperationMap的定义，其中定义支持的op类型，以及op计算公式模板。将子图匹配和代码生成中支持的op类型关联起来。
• 调整了一下代码生成的代码，可以支持生成grad_op的计算代码。
• 调整代码生成单测的代码，添加生成grad_op代码的单测。

[Xreki]

单测里面生成的代码：

extern "C" __global__ void fused_elementwise_0(int N, float* var0, float* var1, float* var3, float* var5, float* var2, float* var4, float* var6, float* var7, float* var8) {
  for(int idx = blockIdx.x * blockDim.x + threadIdx.x;
        idx < N;
     idx += gridDim.x * blockDim.x) {
     var2[idx] = var0[idx] * var1[idx];
     var4[idx] = var2[idx] + var3[idx];
     var6[idx] = var4[idx] - var5[idx];
     var7[idx] = real_max(var6[idx], 0);
     var8[idx] = 1.0 / (1.0 + real_exp(- var7[idx]));
   }
}

extern "C" __global__ void fused_elementwise_grad_0(int N, float* var0, float* var1, float* var2, float* var3, float* var7, float* var4, float* var5, float* var6) {
   for(int idx = blockIdx.x * blockDim.x + threadIdx.x;
       idx < N;
       idx += gridDim.x * blockDim.x) {
     var6[idx] = var2[idx] > 0 ? var7[idx] : 0;
     var4[idx] = var6[idx] * var1[idx];
     var5[idx] = var6[idx] * var0[idx];
   }
}

目前存在的一些问题：

1. fused_elementwise_0中，一些中间计算结果，需要在后面使用到，应该用临时变量保存起来，避免先写入显存，后面再从显存中读取出来，比如var2[idx]、var4[idx]、var6[idx]、var7[idx]。
2. 实际融合的计算中，一些中间结果可能是不需要写入显存的，即我们只需要拿到var8[idx]，因此var2[idx]、var4[idx]、var6[idx]、var7[idx]这些都是中间计算结果，都不需要写入显存。
3. 在梯度计算中，有些grad_op计算需要用到输入x、y和out，而有些不需要用到。在定义每个operation的计算模板的时候，为了明确每个参数的函数，所有输入都考虑进去了。在生成代码时，需要能够过滤掉实际不需要访问到的变量，比如fused_elementwise_grad_0中的float* var3。

[wangchaochaohu]

这里为啥num_operands 小于0

[Xreki]

Find函数用于获取OperationMap内指定type、num_operands的op_type。当num_operands < 0，或者可以是== 0时，代表获取指定type的所有op_type，包括任意操作数的。

[wangchaochaohu]

这里面的expr的size为何不能为1

[Xreki]

目前的设定是，所有的Operation，前向计算只有一个输出，可以有多个输入。所以如果是前向算子，则应该只有1个计算公式。这里判断，如果这个算是不是反向算子（是前向算子），且不是1个计算公式，则认为是不合法的。

[wangchaochaohu]

LGTM