From e92d879a91fab88537355bf142c987c75adc59b0 Mon Sep 17 00:00:00 2001 From: Hu Shenwei Date: Wed, 13 Aug 2025 17:54:47 +0800 Subject: [PATCH 1/5] fix(activation_kernel.cc/cu): fix output type diff for floor/ceil kernel --- .../generator/eager_gen.py | 7 +---- .../fluid/pir/dialect/op_generator/api_gen.py | 7 +---- .../phi/kernels/cpu/activation_grad_kernel.cc | 30 +++++++++++++++++-- paddle/phi/kernels/cpu/activation_kernel.cc | 30 +++++++++++++++++-- .../phi/kernels/gpu/activation_grad_kernel.cu | 28 +++++++++++++++-- paddle/phi/kernels/gpu/activation_kernel.cu | 28 +++++++++++++++-- paddle/phi/kernels/xpu/activation_kernel.cc | 11 ++++++- python/paddle/tensor/ops.py | 1 - 8 files changed, 120 insertions(+), 22 deletions(-) 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 ee95ac3da7d3a7..ca2e409add71d0 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 f266e480b172d2..a7e6c81e5d13da 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 1d98dba03d9993..f91e5a77a57149 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 624c25a94e63be..f110433773303d 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/gpu/activation_grad_kernel.cu b/paddle/phi/kernels/gpu/activation_grad_kernel.cu index 54193a478e9a0e..55b4ae0fd1f1cd 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 718cb08e3013e6..3e2e87527d61ed 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 9913f3eb7f7e2b..89713483efcd2f 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/python/paddle/tensor/ops.py b/python/paddle/tensor/ops.py index 8a29e586241a7b..642e2380fa749d 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 From 8f39507f67bd561318ce0c52c6068238a9b96007 Mon Sep 17 00:00:00 2001 From: Hu Shenwei Date: Fri, 15 Aug 2025 14:47:18 +0800 Subject: [PATCH 2/5] fix(test_activation_op.py): add unit test --- test/legacy_test/test_activation_op.py | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/test/legacy_test/test_activation_op.py b/test/legacy_test/test_activation_op.py index a40ce6f718094d..ce32beebdedc5a 100644 --- a/test/legacy_test/test_activation_op.py +++ b/test/legacy_test/test_activation_op.py @@ -2077,6 +2077,11 @@ def init_shape(self): self.shape = [] +class TestCeil_Int16(TestCeil): + def init_dtype(self): + self.dtype = np.int16 + + class TestFloor(TestActivation): def setUp(self): self.op_type = "floor" @@ -2138,6 +2143,11 @@ def init_shape(self): self.shape = [] +class TestFloor_Int16(TestFloor): + def init_dtype(self): + self.dtype = np.int16 + + class TestCos(TestActivation): def setUp(self): self.op_type = "cos" From ba0e9fb2de5c1a9ad18f996555a1729a6ceba903 Mon Sep 17 00:00:00 2001 From: Hu Shenwei Date: Fri, 15 Aug 2025 17:49:10 +0800 Subject: [PATCH 3/5] fix(full_kernel.cc/cu): add int8 support for full_like --- paddle/phi/kernels/cpu/full_kernel.cc | 1 + paddle/phi/kernels/xpu/full_kernel.cc | 1 + 2 files changed, 2 insertions(+) diff --git a/paddle/phi/kernels/cpu/full_kernel.cc b/paddle/phi/kernels/cpu/full_kernel.cc index c1b0d7de00bf13..188ee2e7bb09fe 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/xpu/full_kernel.cc b/paddle/phi/kernels/xpu/full_kernel.cc index 52706afa7c806b..1854431c3c0a35 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, From e5bfaa641652b1364c0818a5baf2a53e18d01f44 Mon Sep 17 00:00:00 2001 From: Hu Shenwei Date: Sun, 17 Aug 2025 00:36:59 +0800 Subject: [PATCH 4/5] fix(activation_functor.h): fix floor/ceil functor for int dtype input --- paddle/phi/kernels/funcs/activation_functor.h | 44 +++++++++++++++++-- 1 file changed, 40 insertions(+), 4 deletions(-) diff --git a/paddle/phi/kernels/funcs/activation_functor.h b/paddle/phi/kernels/funcs/activation_functor.h index fa55cd725f8319..4d516663ae302d 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)); + } } }; From fbbe8b255f0913e24d920b2ce9a641847ba9a8a4 Mon Sep 17 00:00:00 2001 From: Hu Shenwei Date: Sun, 17 Aug 2025 19:36:20 +0800 Subject: [PATCH 5/5] fix(test_activation_op.py): add unit test --- test/legacy_test/test_activation_op.py | 40 ++++++++++++++++++++++++++ 1 file changed, 40 insertions(+) diff --git a/test/legacy_test/test_activation_op.py b/test/legacy_test/test_activation_op.py index ce32beebdedc5a..167aec86f1a469 100644 --- a/test/legacy_test/test_activation_op.py +++ b/test/legacy_test/test_activation_op.py @@ -2077,11 +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" @@ -2143,11 +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"