diff --git a/paddle/fluid/eager/auto_code_generator/generator/eager_gen.py b/paddle/fluid/eager/auto_code_generator/generator/eager_gen.py index ee95ac3da7d3a..ca2e409add71d 100644 --- a/paddle/fluid/eager/auto_code_generator/generator/eager_gen.py +++ b/paddle/fluid/eager/auto_code_generator/generator/eager_gen.py @@ -152,13 +152,11 @@ "asinh": ["x"], "atan": ["x"], "atanh": ["x"], - "ceil": ["x"], "cos": ["x"], "cosh": ["x"], "digamma": ["x"], "erf": ["x"], "erfinv": ["x"], - "floor": ["x"], "i0": ["x"], "i0e": ["x"], "i1": ["x"], @@ -181,10 +179,7 @@ # ops support casting int tensor into float32 to do forward calculation, # and it is valid to cast float32 gradient back to int tensor. -type_autocast_valid_grad_op_list = { - "ceil", - "floor", -} +type_autocast_valid_grad_op_list = {} # dict of special api that forward api's output will affect backward api's output # backward api's output usually affected by backward api's input diff --git a/paddle/fluid/pir/dialect/op_generator/api_gen.py b/paddle/fluid/pir/dialect/op_generator/api_gen.py index f266e480b172d..a7e6c81e5d13d 100644 --- a/paddle/fluid/pir/dialect/op_generator/api_gen.py +++ b/paddle/fluid/pir/dialect/op_generator/api_gen.py @@ -83,13 +83,11 @@ "asinh": ["x"], "atan": ["x"], "atanh": ["x"], - "ceil": ["x"], "cos": ["x"], "cosh": ["x"], "digamma": ["x"], "erf": ["x"], "erfinv": ["x"], - "floor": ["x"], "i0": ["x"], "i0e": ["x"], "i1": ["x"], @@ -112,10 +110,7 @@ # ops support casting int tensor into float32 to do forward calculation, # and it is valid to cast float32 gradient back to int tensor. -type_autocast_valid_grad_op_list = { - "ceil", - "floor", -} +type_autocast_valid_grad_op_list = {} PD_MANUAL_API_LIST = { 'embedding_grad', diff --git a/paddle/phi/kernels/cpu/activation_grad_kernel.cc b/paddle/phi/kernels/cpu/activation_grad_kernel.cc index 1d98dba03d999..f91e5a77a5714 100644 --- a/paddle/phi/kernels/cpu/activation_grad_kernel.cc +++ b/paddle/phi/kernels/cpu/activation_grad_kernel.cc @@ -483,8 +483,6 @@ PD_REGISTER_ACTIVATION_DOUBLE_GRAD_KERNEL_WITH_COMPLEX(log_double_grad, PD_REGISTER_ACTIVATION_GRAD_KERNEL_WITH_COMPLEX(hardswish_grad, HardSwishGradKernel) PD_REGISTER_ACTIVATION_GRAD_KERNEL(swish_grad, SwishGradKernel) -PD_REGISTER_ACTIVATION_GRAD_KERNEL(floor_grad, FloorGradKernel) -PD_REGISTER_ACTIVATION_GRAD_KERNEL(ceil_grad, CeilGradKernel) PD_REGISTER_ACTIVATION_GRAD_KERNEL(celu_grad, CeluGradKernel) PD_REGISTER_ACTIVATION_DOUBLE_GRAD_KERNEL(celu_double_grad, CeluDoubleGradKernel) @@ -541,3 +539,31 @@ PD_REGISTER_KERNEL(pow_triple_grad, int64_t, phi::dtype::complex, phi::dtype::complex) {} + +PD_REGISTER_KERNEL(ceil_grad, + CPU, + ALL_LAYOUT, + phi::CeilGradKernel, + float, + double, + uint8_t, + int8_t, + int16_t, + int, + int64_t, + phi::dtype::float16, + phi::dtype::bfloat16) {} + +PD_REGISTER_KERNEL(floor_grad, + CPU, + ALL_LAYOUT, + phi::FloorGradKernel, + float, + double, + uint8_t, + int8_t, + int16_t, + int, + int64_t, + phi::dtype::float16, + phi::dtype::bfloat16) {} diff --git a/paddle/phi/kernels/cpu/activation_kernel.cc b/paddle/phi/kernels/cpu/activation_kernel.cc index 624c25a94e63b..f110433773303 100644 --- a/paddle/phi/kernels/cpu/activation_kernel.cc +++ b/paddle/phi/kernels/cpu/activation_kernel.cc @@ -254,8 +254,6 @@ PD_REGISTER_ACTIVATION_KERNEL(hardsigmoid, HardSigmoidKernel) PD_REGISTER_ACTIVATION_KERNEL(swish, SwishKernel) PD_REGISTER_ACTIVATION_KERNEL(relu6, Relu6Kernel) PD_REGISTER_ACTIVATION_KERNEL_WITH_COMPLEX(hardswish, HardSwishKernel) -PD_REGISTER_ACTIVATION_KERNEL(floor, FloorKernel) -PD_REGISTER_ACTIVATION_KERNEL(ceil, CeilKernel) PD_REGISTER_ACTIVATION_KERNEL(celu, CeluKernel) PD_REGISTER_KERNEL( @@ -381,3 +379,31 @@ PD_REGISTER_KERNEL(pow, int64_t, phi::dtype::complex, phi::dtype::complex) {} + +PD_REGISTER_KERNEL(ceil, + CPU, + ALL_LAYOUT, + phi::CeilKernel, + float, + double, + uint8_t, + int8_t, + int16_t, + int, + int64_t, + phi::dtype::float16, + phi::dtype::bfloat16) {} + +PD_REGISTER_KERNEL(floor, + CPU, + ALL_LAYOUT, + phi::FloorKernel, + float, + double, + uint8_t, + int8_t, + int16_t, + int, + int64_t, + phi::dtype::float16, + phi::dtype::bfloat16) {} diff --git a/paddle/phi/kernels/cpu/full_kernel.cc b/paddle/phi/kernels/cpu/full_kernel.cc index c1b0d7de00bf1..188ee2e7bb09f 100644 --- a/paddle/phi/kernels/cpu/full_kernel.cc +++ b/paddle/phi/kernels/cpu/full_kernel.cc @@ -142,6 +142,7 @@ PD_REGISTER_KERNEL(full_like, float, double, uint8_t, + int8_t, int16_t, int, int64_t, diff --git a/paddle/phi/kernels/funcs/activation_functor.h b/paddle/phi/kernels/funcs/activation_functor.h index fa55cd725f831..4d516663ae302 100644 --- a/paddle/phi/kernels/funcs/activation_functor.h +++ b/paddle/phi/kernels/funcs/activation_functor.h @@ -3036,7 +3036,16 @@ template struct FloorFunctor : public BaseActivationFunctor { template void operator()(Device d, X x, Out out) const { - out.device(d) = x.floor(); + if constexpr ((std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value)) { + out.device(d) = x; + } else { + out.device(d) = x.floor(); + } } }; @@ -3160,7 +3169,16 @@ template struct CeilFunctor : public BaseActivationFunctor { template void operator()(Device d, X x, Out out) const { - out.device(d) = x.ceil(); + if constexpr ((std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value)) { + out.device(d) = x; + } else { + out.device(d) = x.ceil(); + } } }; @@ -5403,7 +5421,16 @@ struct CudaCeilFunctor : public BaseActivationFunctor { // ceil(x) = ceil(x) __device__ __forceinline__ T operator()(const T arg_x) const { MPType x = static_cast(arg_x); - return static_cast(ceil(x)); + if constexpr ((std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value)) { + return static_cast(x); + } else { + return static_cast(ceil(x)); + } } }; @@ -5492,7 +5519,16 @@ struct CudaFloorFunctor : public BaseActivationFunctor { // floor(x) = floor(x) __device__ __forceinline__ T operator()(const T arg_x) const { MPType x = static_cast(arg_x); - return static_cast(floor(x)); + if constexpr ((std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value) || + (std::is_same::value)) { + return static_cast(x); + } else { + return static_cast(floor(x)); + } } }; diff --git a/paddle/phi/kernels/gpu/activation_grad_kernel.cu b/paddle/phi/kernels/gpu/activation_grad_kernel.cu index 54193a478e9a0..55b4ae0fd1f1c 100644 --- a/paddle/phi/kernels/gpu/activation_grad_kernel.cu +++ b/paddle/phi/kernels/gpu/activation_grad_kernel.cu @@ -554,8 +554,6 @@ PD_REGISTER_KERNEL(log_double_grad, PD_REGISTER_ACTIVATION_GRAD_KERNEL_WITH_COMPLEX(hardswish_grad, HardSwishGradKernel) PD_REGISTER_ACTIVATION_GRAD_KERNEL(swish_grad, SwishGradKernel) -PD_REGISTER_ACTIVATION_GRAD_KERNEL(floor_grad, FloorGradKernel) -PD_REGISTER_ACTIVATION_GRAD_KERNEL(ceil_grad, CeilGradKernel) PD_REGISTER_ACTIVATION_GRAD_KERNEL(celu_grad, CeluGradKernel) PD_REGISTER_ACTIVATION_GRAD_KERNEL(celu_double_grad, CeluDoubleGradKernel) @@ -617,3 +615,29 @@ PD_REGISTER_KERNEL(pow_triple_grad, phi::dtype::bfloat16, phi::dtype::complex, phi::dtype::complex) {} +PD_REGISTER_KERNEL(ceil_grad, + GPU, + ALL_LAYOUT, + phi::CeilGradKernel, + float, + double, + uint8_t, + int8_t, + int16_t, + int, + int64_t, + phi::dtype::float16, + phi::dtype::bfloat16) {} +PD_REGISTER_KERNEL(floor_grad, + GPU, + ALL_LAYOUT, + phi::FloorGradKernel, + float, + double, + uint8_t, + int8_t, + int16_t, + int, + int64_t, + phi::dtype::float16, + phi::dtype::bfloat16) {} diff --git a/paddle/phi/kernels/gpu/activation_kernel.cu b/paddle/phi/kernels/gpu/activation_kernel.cu index 718cb08e3013e..3e2e87527d61e 100644 --- a/paddle/phi/kernels/gpu/activation_kernel.cu +++ b/paddle/phi/kernels/gpu/activation_kernel.cu @@ -347,8 +347,6 @@ PD_REGISTER_ACTIVATION_KERNEL_WITH_COMPLEX(logsigmoid, LogSigmoidKernel) PD_REGISTER_ACTIVATION_KERNEL(hardsigmoid, HardSigmoidKernel) PD_REGISTER_ACTIVATION_KERNEL_WITH_COMPLEX(hardswish, HardSwishKernel) PD_REGISTER_ACTIVATION_KERNEL(swish, SwishKernel) -PD_REGISTER_ACTIVATION_KERNEL(floor, FloorKernel) -PD_REGISTER_ACTIVATION_KERNEL(ceil, CeilKernel) PD_REGISTER_ACTIVATION_KERNEL(celu, CeluKernel) PD_REGISTER_ACTIVATION_KERNEL(selu, SeluKernel) PD_REGISTER_ACTIVATION_KERNEL(logit, LogitCUDAKernel) @@ -435,3 +433,29 @@ PD_REGISTER_KERNEL(pow, phi::dtype::bfloat16, phi::dtype::complex, phi::dtype::complex) {} +PD_REGISTER_KERNEL(ceil, + GPU, + ALL_LAYOUT, + phi::CeilKernel, + float, + double, + uint8_t, + int8_t, + int16_t, + int, + int64_t, + phi::dtype::float16, + phi::dtype::bfloat16) {} +PD_REGISTER_KERNEL(floor, + GPU, + ALL_LAYOUT, + phi::FloorKernel, + float, + double, + uint8_t, + int8_t, + int16_t, + int, + int64_t, + phi::dtype::float16, + phi::dtype::bfloat16) {} diff --git a/paddle/phi/kernels/xpu/activation_kernel.cc b/paddle/phi/kernels/xpu/activation_kernel.cc index 9913f3eb7f7e2..89713483efcd2 100644 --- a/paddle/phi/kernels/xpu/activation_kernel.cc +++ b/paddle/phi/kernels/xpu/activation_kernel.cc @@ -777,7 +777,16 @@ PD_REGISTER_KERNEL(acos, #define PD_REGISTER_ACTIVATION_KERNEL(name, func) \ PD_REGISTER_KERNEL(name, XPU, ALL_LAYOUT, phi::func, float) {} -PD_REGISTER_ACTIVATION_KERNEL(floor, FloorKernel) PD_REGISTER_ACTIVATION_KERNEL(mish, MishKernel) PD_REGISTER_ACTIVATION_KERNEL(reciprocal, ReciprocalKernel) PD_REGISTER_ACTIVATION_KERNEL(softplus, SoftplusKernel) + +PD_REGISTER_KERNEL(floor, + XPU, + ALL_LAYOUT, + phi::FloorKernel, + float, + int, + int64_t, + phi::dtype::float16, + phi::dtype::bfloat16) {} diff --git a/paddle/phi/kernels/xpu/full_kernel.cc b/paddle/phi/kernels/xpu/full_kernel.cc index 52706afa7c806..1854431c3c0a3 100644 --- a/paddle/phi/kernels/xpu/full_kernel.cc +++ b/paddle/phi/kernels/xpu/full_kernel.cc @@ -174,6 +174,7 @@ PD_REGISTER_KERNEL(full_like, float, double, uint8_t, + int8_t, int16_t, int, int64_t, diff --git a/python/paddle/tensor/ops.py b/python/paddle/tensor/ops.py index 8a29e586241a7..642e2380fa749 100644 --- a/python/paddle/tensor/ops.py +++ b/python/paddle/tensor/ops.py @@ -701,7 +701,6 @@ def floor(x: Tensor, name: str | None = None) -> Tensor: Returns: Tensor. Output of Floor operator, a Tensor with shape same as input - (integer types are autocasted into float32). Examples: .. code-block:: python diff --git a/test/legacy_test/test_activation_op.py b/test/legacy_test/test_activation_op.py index a40ce6f718094..167aec86f1a46 100644 --- a/test/legacy_test/test_activation_op.py +++ b/test/legacy_test/test_activation_op.py @@ -2077,6 +2077,31 @@ def init_shape(self): self.shape = [] +class TestCeil_UInt8(TestCeil): + def init_dtype(self): + self.dtype = np.uint8 + + +class TestCeil_Int8(TestCeil): + def init_dtype(self): + self.dtype = np.int8 + + +class TestCeil_Int16(TestCeil): + def init_dtype(self): + self.dtype = np.int16 + + +class TestCeil_Int32(TestCeil): + def init_dtype(self): + self.dtype = np.int32 + + +class TestCeil_Int64(TestCeil): + def init_dtype(self): + self.dtype = np.int64 + + class TestFloor(TestActivation): def setUp(self): self.op_type = "floor" @@ -2138,6 +2163,31 @@ def init_shape(self): self.shape = [] +class TestFloor_UInt8(TestFloor): + def init_dtype(self): + self.dtype = np.uint8 + + +class TestFloor_Int8(TestFloor): + def init_dtype(self): + self.dtype = np.int8 + + +class TestFloor_Int16(TestFloor): + def init_dtype(self): + self.dtype = np.int16 + + +class TestFloor_Int32(TestFloor): + def init_dtype(self): + self.dtype = np.int32 + + +class TestFloor_Int64(TestFloor): + def init_dtype(self): + self.dtype = np.int64 + + class TestCos(TestActivation): def setUp(self): self.op_type = "cos"