| #include "caffe2/operators/cos_op.h" | 
 |  | 
 | #include <algorithm> | 
 | #include <functional> | 
 |  | 
 | #include "caffe2/core/context_gpu.h" | 
 |  | 
 | namespace caffe2 { | 
 |  | 
 | namespace { | 
 |  | 
 | template <typename T> | 
 | __global__ void | 
 | CosGradientCUDAKernel(const int N, const T* dY, const T* X, T* dX) { | 
 |   CUDA_1D_KERNEL_LOOP(i, N) { | 
 | #if __CUDA_ARCH__ >= 350 | 
 |     dX[i] = -__ldg(dY + i) * sin(__ldg(X + i)); | 
 | #else | 
 |     dX[i] = -dY[i] * sin(X[i]); | 
 | #endif | 
 |   } | 
 | } | 
 |  | 
 | } // namespace | 
 |  | 
 | template <> | 
 | template <typename T> | 
 | bool CosGradientFunctor<CUDAContext>::Forward( | 
 |     const std::vector<int>& X_dims, | 
 |     const std::vector<int>& /* dY_dims */, | 
 |     const T* X, | 
 |     const T* dY, | 
 |     T* dX, | 
 |     CUDAContext* context) const { | 
 |   const int size = std::accumulate( | 
 |       X_dims.cbegin(), X_dims.cend(), 1, std::multiplies<int>()); | 
 |   CosGradientCUDAKernel<<< | 
 |       CAFFE_GET_BLOCKS(size), | 
 |       CAFFE_CUDA_NUM_THREADS, | 
 |       0, | 
 |       context->cuda_stream()>>>(size, dY, X, dX); | 
 |   return true; | 
 | } | 
 |  | 
 | REGISTER_CUDA_OPERATOR( | 
 |     Cos, | 
 |     UnaryElementwiseOp< | 
 |         TensorTypes<float>, | 
 |         CUDAContext, | 
 |         CosFunctor<CUDAContext>>); | 
 | REGISTER_CUDA_OPERATOR( | 
 |     CosGradient, | 
 |     BinaryElementwiseOp< | 
 |         TensorTypes<float>, | 
 |         CUDAContext, | 
 |         CosGradientFunctor<CUDAContext>>); | 
 |  | 
 | } // namespace caffe2 |