Implement GPU kernel for cross entropy operator.

[qingqing01]

Fix #3492

• Implement GPU kernel for cross entropy operator.
• Fix the origin CPU implementation.

[jacquesqiao]

why not just use a function

template<typename T> 
__host__ __device__ T clipping_log(const Tx) {
  ...
  return xx;
}

[qingqing01]

I notice that Majel and WarpCTC use many functors to define many common functions. The functor has its own advantages, see: https://stackoverflow.com/questions/356950/c-functors-and-their-uses

But it is not necessary here, I change it to a plain function.

[QiJune]

我有点没搞清楚啊，交叉熵的计算公式是什么呀？不是 cost = 求和(yln(x) + (1-y)ln(1-x)) 这个吗？

[qingqing01]

label实际是one-hot类型: [0, 0, ..., 1, ..., 0,0] ，所以这里是：y[i] = -log(x[ i * D + label[i] ]), 求和通常不在cross_entropy_op里做，是后面接的op或者打印loss时再求和。

[zchen0211]

Hello, do we need something like <<<grid, block, 0, stream_id >>> to execute it?

[qingqing01]

@zchen0211 We need to specify the stream, like <<<grid, block, 0, stream_id >>>. I think the stream_id can be got from CUDAContext, but there is no stream in CUDAContext now. And @QiJune said he has worked on it last week.

[QiJune]

LGTM