template<typename Dtype>inline int8_t caffe_sign(Dtype val) { return (Dtype(0) < val) - (val < Dtype(0));}
// The following two macros are modifications of DEFINE_VSL_UNARY_FUNC// in include/caffe/util/mkl_alternate.hpp authored by @Rowland Depp.// Please refer to commit 7e8ef25c7 of the boost-eigen branch.// Git cherry picking that commit caused a conflict hard to resolve and// copying that file in convenient for code reviewing.// So they have to be pasted here temporarily.#define DEFINE_CAFFE_CPU_UNARY_FUNC(name, operation) \ template<typename Dtype> \ void caffe_cpu_##name(const int n, const Dtype* x, Dtype* y) { \ CHECK_GT(n, 0); CHECK(x); CHECK(y); \ for (int i = 0; i < n; ++i) { \ operation; \ } \ }
// output is 1 for the positives, 0 for zero, and -1 for the negativesDEFINE_CAFFE_CPU_UNARY_FUNC(sign, y[i] = caffe_sign<Dtype>(x[i]));
#define DEFINE_AND_INSTANTIATE_GPU_UNARY_FUNC(name, operation) \template<typename Dtype> \__global__ void name##_kernel(const int n, const Dtype* x, Dtype* y) { \ CUDA_KERNEL_LOOP(index, n) { \ operation; \ } \} \template <> \void caffe_gpu_##name<float>(const int n, const float* x, float* y) { \ /* NOLINT_NEXT_LINE(whitespace/operators) */ \ name##_kernel<float><<<CAFFE_GET_BLOCKS(n), CAFFE_CUDA_NUM_THREADS>>>( \ n, x, y); \} \template <> \void caffe_gpu_##name<double>(const int n, const double* x, double* y) { \ /* NOLINT_NEXT_LINE(whitespace/operators) */ \ name##_kernel<double><<<CAFFE_GET_BLOCKS(n), CAFFE_CUDA_NUM_THREADS>>>( \ n, x, y); \}
DEFINE_AND_INSTANTIATE_GPU_UNARY_FUNC(sign, y[index] = (Dtype(0) < x[index]) - (x[index] < Dtype(0)));