text: /************************************************************
   *
   * Licensed to the Apache Software Foundation (ASF) under one
   * or more contributor license agreements.  See the NOTICE file
   * distributed with this work for additional information
   * regarding copyright ownership.  The ASF licenses this file
   * to you under the Apache License, Version 2.0 (the
   * "License"); you may not use this file except in compliance
   * with the License.  You may obtain a copy of the License at
   *
   *   http://www.apache.org/licenses/LICENSE-2.0
   *
   * Unless required by applicable law or agreed to in writing,
   * software distributed under the License is distributed on an
   * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
   * KIND, either express or implied.  See the License for the
   * specific language governing permissions and limitations
   * under the License.
   *
   *************************************************************/
  #ifndef SRC_CORE_TENSOR__MATH_KERNEL_H_
  #define SRC_CORE_TENSOR__MATH_KERNEL_H_
  
  #include "singa/singa_config.h"
  #ifdef USE_CUDA
  
  #include <thrust/execution_policy.h>
  #include <thrust/remove.h>
  #include <thrust/sort.h>
  
  #include "cuda_fp16.h"
  
  /// TODO(wangwei) Clean the function APIs as commented in tensor_math.h
  ///  Add 'Context *ctx' as an argument of all cuda functions.
  namespace singa {
  
  // TODO(wangwei) make all function templates.
  namespace cuda {
  
  // 0 input
  void set(const size_t n, const float v, float *out, cudaStream_t s);
  
  // 1 input
  void abs(const size_t n, const float *in, float *out, cudaStream_t s);
  void sign(const size_t n, const float *in, float *out, cudaStream_t s);
  void exp(const size_t n, const float *in, float *out, cudaStream_t s);
  void erf(const size_t n, const float *in, float *out, cudaStream_t s);
  void ceil2(const size_t n, const float *in, float *out, cudaStream_t s);
  void floor(const size_t n, const float *in, float *out, cudaStream_t s);
  void round(const size_t n, const float *in, float *out, cudaStream_t s);
  void rounde(const size_t n, const float *in, float *out, cudaStream_t s);
  void cast_float_2_int(const size_t n, const float *src, int *dst,
                        cudaStream_t s);
  void cast_int_2_float(const size_t n, const int *src, float *dst,
                        cudaStream_t s);
  void log(const size_t n, const float *in, float *out, cudaStream_t s);
  void sqrt(const size_t n, const float *in, float *out, cudaStream_t s);
  void square(const size_t n, const float *in, float *out, cudaStream_t s);
  void cos(const size_t n, const float *in, float *out, cudaStream_t s);
  void cosh(const size_t n, const float *in, float *out, cudaStream_t s);
  void acos(const size_t n, const float *in, float *out, cudaStream_t s);
  void acosh(const size_t n, const float *in, float *out, cudaStream_t s);
  void sin(const size_t n, const float *in, float *out, cudaStream_t s);
  void sinh(const size_t n, const float *in, float *out, cudaStream_t s);
  void asin(const size_t n, const float *in, float *out, cudaStream_t s);
  void asinh(const size_t n, const float *in, float *out, cudaStream_t s);
  void tan(const size_t n, const float *in, float *out, cudaStream_t s);
  void tanh(const size_t n, const float *in, float *out, cudaStream_t s);
  void atan(const size_t n, const float *in, float *out, cudaStream_t s);
  void atanh(const size_t n, const float *in, float *out, cudaStream_t s);
  void relu(const size_t n, const float *in, float *out, cudaStream_t s);
  void relu(const size_t n, const __half *in, __half *out, cudaStream_t s);
  void sigmoid(const size_t n, const float *in, float *out, cudaStream_t s);
  void softplus(const size_t n, const float *in, float *out, cudaStream_t s);
  void softsign(const size_t n, const float *in, float *out, cudaStream_t s);
  void clamp(const size_t n, const float low, const float high, const float *in,
             float *out, cudaStream_t s);
  
  void pow(const size_t n, const float *in, const float x, float *out,
           cudaStream_t s);
  
  void add(const size_t n, const float *in, const float x, float *out,
           cudaStream_t s);
  
  void mult(const size_t n, const float *in, const float x, float *out,
            cudaStream_t s);
  void mult(const size_t n, const __half *in, const __half x, __half *out,
            cudaStream_t s);
  
  void traverse_unary_transform(const size_t n, size_t nDim, const float *in,
                                const int *shape, const int *stride, float *out,
                                cudaStream_t s);
  void traverse_unary_transform(const size_t n, size_t nDim, const __half *in,
                                const int *shape, const int *stride, __half *out,
                                cudaStream_t s);
  
  void div(const size_t n, const float x, const float *in, float *out,
           cudaStream_t s);
  
  void threshold(const size_t n, const float x, const float *in, float *out,
                 cudaStream_t s);
  
  void relubackward(const size_t num, const float *in1, const float *in2,
                    float *out, cudaStream_t s);
  
  void relubackward(const size_t num, const __half *in1, const __half *in2,
                    __half *out, cudaStream_t s);
  
  void gt(const size_t num, const float *in, const float x, float *out,
          cudaStream_t s);
  void gt(const size_t num, const float *in1, const float *in2, float *out,
          cudaStream_t s);
  
  void ge(const size_t num, const float *in, const float x, float *out,
          cudaStream_t s);
  void ge(const size_t num, const float *in1, const float *in2, float *out,
          cudaStream_t s);
  
  void eq(const size_t num, const float *in, const float x, float *out,
          cudaStream_t s);
  void eq(const size_t num, const float *in1, const float *in2, float *out,
          cudaStream_t s);
  
  void lt(const size_t num, const float *in, const float x, float *out,
          cudaStream_t s);
  void lt(const size_t num, const float *in1, const float *in2, float *out,
          cudaStream_t s);
  
  void le(const size_t num, const float *in, const float x, float *out,
          cudaStream_t s);
  void le(const size_t num, const float *in1, const float *in2, float *out,
          cudaStream_t s);
  
  // 2 inputs
  void pow(const size_t n, const float *in1, const float *in2, float *out,
           cudaStream_t s);
  void pow(const size_t n, const __half *in1, const __half *in2, __half *out,
           cudaStream_t s);
  
  void add(const size_t n, const float *in1, const float *in2, float *out,
           cudaStream_t s);
  void add(const size_t n, const __half *in1, const __half *in2, __half *out,
           cudaStream_t s);
  
  void sub(const size_t n, const float *in1, const float *in2, float *out,
           cudaStream_t s);
  void sub(const size_t n, const __half *in1, const __half *in2, __half *out,
           cudaStream_t s);
  
  void mult(const size_t n, const float *in1, const float *in2, float *out,
            cudaStream_t s);
  void mult(const size_t n, const __half *in1, const __half *in2, __half *out,
            cudaStream_t s);
  
  void div(const size_t n, const float *in1, const float *in2, float *out,
           cudaStream_t s);
  void div(const size_t n, const __half *in1, const __half *in2, __half *out,
           cudaStream_t s);
  
  // void sum(const size_t n, const float *in, float *out, cudaStream_t s);
  
  void ComputeCrossEntropy(bool int_target, const size_t batchsize,
                           const size_t dim, const float *p, const int *t,
                           float *loss, cudaStream_t stream);
  void ComputeCrossEntropy(bool int_target, const size_t batchsize,
                           const size_t dim, const __half *p, const int *t,
                           __half *loss, cudaStream_t stream);
  void SoftmaxCrossEntropyBwd(bool int_target, const size_t batchsize,
                              const size_t dim, const float *p, const int *t,
                              float *grad, cudaStream_t stream);
  void SoftmaxCrossEntropyBwd(bool int_target, const size_t batchsize,
                              const size_t dim, const __half *p, const int *t,
                              __half *grad, cudaStream_t stream);
  
  void RowMax(const size_t nrow, const size_t ncol, const float *inPtr,
              float *outPtr, cudaStream_t stream);
  
  void float2half(const size_t n, const float *in, __half *out, cudaStream_t s);
  
  void half2float(const size_t n, const __half *in, float *out, cudaStream_t s);
  
  void sparsabs(const size_t n, const float threshold, const float *in,
                float *out, cudaStream_t s);
  
  void sparsindex(const size_t n, const float *in, int *out, cudaStream_t s);
  
  void generateindex(const size_t n, int *out, cudaStream_t s);
  
  void removezeroval(const size_t n, float *in, cudaStream_t s);
  
  void removezeroidx(const size_t n, int *in, cudaStream_t s, int *address);
  
  void sortbykey(const size_t n, float *key, int *value, cudaStream_t s);
  
  }  // namespace cuda
  
  }  // namespace singa
  
  #endif
  #endif  // SRC_CORE_TENSOR__MATH_KERNEL_H_
  