/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved. Licensed 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. */ #include #include "paddle/fluid/operators/controlflow/compare_all_op.h" #include "paddle/fluid/operators/elementwise/elementwise_op_impl.cu.h" #include "paddle/fluid/operators/reduce_ops/cub_reduce.h" namespace ops = paddle::operators; namespace plat = paddle::platform; namespace paddle { namespace operators { template struct IdentityFunctor { HOSTDEVICE explicit inline IdentityFunctor() {} HOSTDEVICE inline T operator()(const T& x) const { return x; } }; struct BitwiseAdd { // Bitwise add operator, returns a + b template __host__ __device__ __forceinline__ T operator()(const T& a, const T& b) const { return a & b; } }; template struct CudaEqualReduceFunctor { using ELEM_TYPE = T; HOSTDEVICE bool operator()(const T args[]) const { return (args[0] == args[1]); } }; template struct CudaEqualReduceFunctor< T, typename std::enable_if::value>::type> { using ELEM_TYPE = T; HOSTDEVICE bool operator()(const T args[]) const { return fabs(static_cast(args[0] - args[1])) < 1e-8; } }; template class CompareReduceOpKernel : public framework::OpKernel { public: void Compute(const framework::ExecutionContext& context) const override { using T = typename Functor::ELEM_TYPE; using Tensor = framework::Tensor; auto* x = context.Input("X"); auto* y = context.Input("Y"); auto* z = context.Output("Out"); bool* z_data = z->mutable_data(context.GetPlace()); Tensor tmp; if (x->dims() != y->dims()) { thrust::device_ptr z_dev_ptr(z_data); thrust::fill(z_dev_ptr, z_dev_ptr + 1, false); return; } else { tmp.mutable_data(x->dims(), context.GetPlace()); const auto& cuda_ctx = context.template device_context(); std::vector ins = {x, y}; std::vector outs = {&tmp}; LaunchSameDimsElementwiseCudaKernel( cuda_ctx, ins, &outs, Functor()); // Reduce by 'bitwise and' operator std::vector reduce_dims; reduce_dims.resize(tmp.dims().size()); for (int i = 0; i < reduce_dims.size(); ++i) reduce_dims[i] = i; auto stream = context.cuda_device_context().stream(); TensorReduce>( tmp, z, reduce_dims, true, BitwiseAdd(), IdentityFunctor(), stream); } } }; } // namespace operators } // namespace paddle #define REGISTER_COMPARE_REDUCE_CUDA_KERNEL(op_type, functor) \ REGISTER_OP_CUDA_KERNEL( \ op_type, \ ops::CompareReduceOpKernel>, \ ops::CompareReduceOpKernel>, \ ops::CompareReduceOpKernel>, \ ops::CompareReduceOpKernel>, \ ops::CompareReduceOpKernel>); REGISTER_COMPARE_REDUCE_CUDA_KERNEL(equal_all, CudaEqualReduceFunctor) #undef REGISTER_COMPARE_REDUCE_CUDA_KERNEL