From 742ea80c0b028c3fd4e582dc4a5409e2e82863e5 Mon Sep 17 00:00:00 2001 From: Wang Huan Date: Tue, 14 May 2024 02:59:53 +0000 Subject: [PATCH 1/6] fix pir api 1 --- test/deprecated/legacy_test/test_pool2d_op.py | 1326 ----------------- 1 file changed, 1326 deletions(-) delete mode 100644 test/deprecated/legacy_test/test_pool2d_op.py diff --git a/test/deprecated/legacy_test/test_pool2d_op.py b/test/deprecated/legacy_test/test_pool2d_op.py deleted file mode 100644 index b2f10e3af1b266..00000000000000 --- a/test/deprecated/legacy_test/test_pool2d_op.py +++ /dev/null @@ -1,1326 +0,0 @@ -# Copyright (c) 2018 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. - -import unittest - -import numpy as np -from op_test import OpTest, convert_float_to_uint16 - -import paddle -from paddle.base import core -from paddle.framework import in_dynamic_mode - - -def adaptive_start_index(index, input_size, output_size): - return int(np.floor(index * input_size / output_size)) - - -def adaptive_end_index(index, input_size, output_size): - return int(np.ceil((index + 1) * input_size / output_size)) - - -def max_pool2D_forward_naive( - x, - ksize, - strides, - paddings, - global_pool=0, - ceil_mode=False, - exclusive=True, - adaptive=False, - data_type=np.float64, -): - if data_type == np.float64 and core.is_compiled_with_rocm(): - data_type = np.float32 - N, C, H, W = x.shape - if global_pool == 1: - ksize = [H, W] - if adaptive: - H_out, W_out = ksize - else: - H_out = ( - (H - ksize[0] + 2 * paddings[0] + strides[0] - 1) // strides[0] + 1 - if ceil_mode - else (H - ksize[0] + 2 * paddings[0]) // strides[0] + 1 - ) - W_out = ( - (W - ksize[1] + 2 * paddings[1] + strides[1] - 1) // strides[1] + 1 - if ceil_mode - else (W - ksize[1] + 2 * paddings[1]) // strides[1] + 1 - ) - out = np.zeros((N, C, H_out, W_out)) - for i in range(H_out): - for j in range(W_out): - if adaptive: - r_start = adaptive_start_index(i, H, ksize[0]) - r_end = adaptive_end_index(i, H, ksize[0]) - c_start = adaptive_start_index(j, W, ksize[1]) - c_end = adaptive_end_index(j, W, ksize[1]) - else: - r_start = np.max((i * strides[0] - paddings[0], 0)) - r_end = np.min((i * strides[0] + ksize[0] - paddings[0], H)) - c_start = np.max((j * strides[1] - paddings[1], 0)) - c_end = np.min((j * strides[1] + ksize[1] - paddings[1], W)) - x_masked = x[:, :, r_start:r_end, c_start:c_end] - - out[:, :, i, j] = np.max(x_masked, axis=(2, 3)) - return out - - -def avg_pool2D_forward_naive( - x, - ksize, - strides, - paddings, - global_pool=0, - ceil_mode=False, - exclusive=True, - adaptive=False, - data_type=np.float64, -): - if data_type == np.float64 and core.is_compiled_with_rocm(): - data_type = np.float32 - N, C, H, W = x.shape - if global_pool == 1: - ksize = [H, W] - if adaptive: - H_out, W_out = ksize - else: - H_out = ( - (H - ksize[0] + 2 * paddings[0] + strides[0] - 1) // strides[0] + 1 - if ceil_mode - else (H - ksize[0] + 2 * paddings[0]) // strides[0] + 1 - ) - W_out = ( - (W - ksize[1] + 2 * paddings[1] + strides[1] - 1) // strides[1] + 1 - if ceil_mode - else (W - ksize[1] + 2 * paddings[1]) // strides[1] + 1 - ) - out = np.zeros((N, C, H_out, W_out)) - for i in range(H_out): - for j in range(W_out): - if adaptive: - r_start = adaptive_start_index(i, H, ksize[0]) - r_end = adaptive_end_index(i, H, ksize[0]) - c_start = adaptive_start_index(j, W, ksize[1]) - c_end = adaptive_end_index(j, W, ksize[1]) - else: - r_start = i * strides[0] - paddings[0] - r_end = i * strides[0] + ksize[0] - paddings[0] - c_start = j * strides[1] - paddings[1] - c_end = j * strides[1] + ksize[1] - paddings[1] - field_size = (r_end - r_start) * (c_end - c_start) - r_start = np.max((r_start, 0)) - r_end = np.min((r_end, H)) - c_start = np.max((c_start, 0)) - c_end = np.min((c_end, W)) - - x_masked = x[:, :, r_start:r_end, c_start:c_end] - - if exclusive or adaptive: - field_size = (r_end - r_start) * (c_end - c_start) - - if data_type == np.int8 or data_type == np.uint8: - out[:, :, i, j] = ( - np.rint(np.sum(x_masked, axis=(2, 3)) / field_size) - ).astype(data_type) - else: - out[:, :, i, j] = ( - np.sum(x_masked, axis=(2, 3)) / field_size - ).astype(data_type) - return out - - -def pool2D_forward_naive( - x, - ksize, - strides, - paddings, - global_pool=0, - ceil_mode=False, - exclusive=True, - adaptive=False, - data_format='NCHW', - pool_type="max", - padding_algorithm="EXPLICIT", -): - # update paddings - def _get_padding_with_SAME(input_shape, pool_size, pool_stride): - padding = [] - for input_size, filter_size, stride_size in zip( - input_shape, pool_size, pool_stride - ): - out_size = int((input_size + stride_size - 1) / stride_size) - pad_sum = np.max( - ((out_size - 1) * stride_size + filter_size - input_size, 0) - ) - pad_0 = int(pad_sum / 2) - pad_1 = int(pad_sum - pad_0) - padding.append(pad_0) - padding.append(pad_1) - return padding - - if isinstance(padding_algorithm, str): - padding_algorithm = padding_algorithm.upper() - if padding_algorithm not in ["SAME", "VALID", "EXPLICIT"]: - raise ValueError( - "Unknown Attr(padding_algorithm): '%s'. " - "It can only be 'SAME' or 'VALID'." % str(padding_algorithm) - ) - - if padding_algorithm == "VALID": - paddings = [0, 0, 0, 0] - if ceil_mode is not False: - raise ValueError( - "When Attr(pool_padding) is \"VALID\", Attr(ceil_mode)" - " must be False. " - "Received ceil_mode: True." - ) - elif padding_algorithm == "SAME": - input_data_shape = [] - if data_format == "NCHW": - input_data_shape = x.shape[2:4] - elif data_format == "NHWC": - input_data_shape = x.shape[1:3] - paddings = _get_padding_with_SAME(input_data_shape, ksize, strides) - - assert len(paddings) == 2 or len(paddings) == 4 - is_sys = True if len(paddings) == 2 else False - - N = x.shape[0] - C, H, W = ( - [x.shape[1], x.shape[2], x.shape[3]] - if data_format == 'NCHW' - else [x.shape[3], x.shape[1], x.shape[2]] - ) - - if global_pool == 1: - ksize = [H, W] - paddings = [0 for _ in range(len(paddings))] - - pad_h_up = paddings[0] if is_sys else paddings[0] - pad_h_down = paddings[0] if is_sys else paddings[1] - pad_w_left = paddings[1] if is_sys else paddings[2] - pad_w_right = paddings[1] if is_sys else paddings[3] - - if adaptive: - H_out, W_out = ksize - else: - H_out = ( - (H - ksize[0] + pad_h_up + pad_h_down + strides[0] - 1) - // strides[0] - + 1 - if ceil_mode - else (H - ksize[0] + pad_h_up + pad_h_down) // strides[0] + 1 - ) - W_out = ( - (W - ksize[1] + pad_w_left + pad_w_right + strides[1] - 1) - // strides[1] - + 1 - if ceil_mode - else (W - ksize[1] + pad_w_left + pad_w_right) // strides[1] + 1 - ) - - out = ( - np.zeros((N, C, H_out, W_out)) - if data_format == 'NCHW' - else np.zeros((N, H_out, W_out, C)) - ) - for i in range(H_out): - if adaptive: - in_h_start = adaptive_start_index(i, H, ksize[0]) - in_h_end = adaptive_end_index(i, H, ksize[0]) - else: - in_h_start = np.max((i * strides[0] - pad_h_up, 0)) - in_h_end = np.min((i * strides[0] + ksize[0] - pad_h_up, H)) - - for j in range(W_out): - if adaptive: - in_w_start = adaptive_start_index(j, W, ksize[1]) - in_w_end = adaptive_end_index(j, W, ksize[1]) - else: - in_h_start = i * strides[0] - pad_h_up - in_w_start = j * strides[1] - pad_w_left - in_h_end = i * strides[0] + ksize[0] - pad_h_up - in_w_end = j * strides[1] + ksize[1] - pad_w_left - - field_size = (in_h_end - in_h_start) * (in_w_end - in_w_start) - in_h_start = np.max((in_h_start, 0)) - in_w_start = np.max((in_w_start, 0)) - in_h_end = np.min((in_h_end, H)) - in_w_end = np.min((in_w_end, W)) - - if data_format == 'NCHW': - x_masked = x[:, :, in_h_start:in_h_end, in_w_start:in_w_end] - if pool_type == 'avg': - if exclusive or adaptive: - field_size = (in_h_end - in_h_start) * ( - in_w_end - in_w_start - ) - - # if (exclusive or adaptive) else (ksize[0] * ksize[1]) - out[:, :, i, j] = np.sum(x_masked, axis=(2, 3)) / field_size - elif pool_type == 'max': - out[:, :, i, j] = np.max(x_masked, axis=(2, 3)) - elif data_format == 'NHWC': - x_masked = x[:, in_h_start:in_h_end, in_w_start:in_w_end, :] - if pool_type == 'avg': - if exclusive or adaptive: - field_size = (in_h_end - in_h_start) * ( - in_w_end - in_w_start - ) - out[:, i, j, :] = np.sum(x_masked, axis=(1, 2)) / field_size - elif pool_type == 'max': - out[:, i, j, :] = np.max(x_masked, axis=(1, 2)) - return out - - -def pool2d_wrapper_not_use_cudnn( - X, - ksize=[], - strides=[], - paddings=[], - ceil_mode=False, - exclusive=True, - data_format="NCDHW", - pooling_type="max", - global_pooling=False, - adaptive=False, - padding_algorithm="EXPLICIT", -): - if in_dynamic_mode(): - X = X._use_gpudnn(False) - if data_format == "AnyLayout": - data_format = "NCDHW" - return paddle._C_ops.pool2d( - X, - ksize, - strides, - paddings, - ceil_mode, - exclusive, - data_format, - pooling_type, - global_pooling, - adaptive, - padding_algorithm, - ) - - -def pool2d_wrapper_use_cudnn( - X, - ksize=[], - strides=[], - paddings=[], - ceil_mode=False, - exclusive=True, - data_format="NCDHW", - pooling_type="max", - global_pooling=False, - adaptive=False, - padding_algorithm="EXPLICIT", -): - if data_format == "AnyLayout": - data_format = "NCDHW" - return paddle._C_ops.pool2d( - X, - ksize, - strides, - paddings, - ceil_mode, - exclusive, - data_format, - pooling_type, - global_pooling, - adaptive, - padding_algorithm, - ) - - -class TestPool2D_Op_Mixin: - def setUp(self): - self.op_type = "pool2d" - self.use_cudnn = False - self.init_kernel_type() - self.use_mkldnn = False - self.init_data_type() - self.init_test_case() - self.padding_algorithm = "EXPLICIT" - self.init_paddings() - self.init_global_pool() - self.init_kernel_type() - self.init_pool_type() - self.init_ceil_mode() - self.init_exclusive() - self.init_adaptive() - self.init_data_format() - self.init_shape() - - if self.is_bfloat16_op(): - input = np.random.random(self.shape).astype(np.float32) - else: - input = np.random.random(self.shape).astype(self.dtype) - - output = pool2D_forward_naive( - input, - self.ksize, - self.strides, - self.paddings, - self.global_pool, - self.ceil_mode, - self.exclusive, - self.adaptive, - self.data_format, - self.pool_type, - self.padding_algorithm, - ) - - if self.is_bfloat16_op(): - output = convert_float_to_uint16(output) - self.inputs = {'X': convert_float_to_uint16(input)} - else: - output = output.astype(self.dtype) - self.inputs = {'X': OpTest.np_dtype_to_base_dtype(input)} - - self.attrs = { - 'strides': self.strides, - 'paddings': self.paddings, - 'ksize': self.ksize, - 'pooling_type': self.pool_type, - 'global_pooling': self.global_pool, - 'use_cudnn': self.use_cudnn, - 'use_mkldnn': self.use_mkldnn, - 'ceil_mode': self.ceil_mode, - 'data_format': self.data_format, - 'exclusive': self.exclusive, - 'adaptive': self.adaptive, - "padding_algorithm": self.padding_algorithm, - } - - self.outputs = {'Out': output} - - if self.use_cudnn: - self.python_api = pool2d_wrapper_use_cudnn - else: - self.python_api = pool2d_wrapper_not_use_cudnn - - def has_cudnn(self): - return core.is_compiled_with_cuda() and self.use_cudnn - - def test_check_output(self): - # TODO(wangzhongpu): support onednn op in dygraph mode - if self.has_cudnn(): - place = core.CUDAPlace(0) - self.check_output_with_place( - place, - atol=1e-5, - check_dygraph=(not self.use_mkldnn), - check_cinn=True, - check_pir=True, - check_pir_onednn=self.check_pir_onednn, - ) - else: - self.check_output( - check_dygraph=(not self.use_mkldnn), - check_pir=True, - check_pir_onednn=self.check_pir_onednn, - ) - - def test_check_grad(self): - if self.dtype == np.float16: - return - # TODO(wangzhongpu): support onednn op in dygraph mode - if self.has_cudnn() and self.pool_type != "max": - place = core.CUDAPlace(0) - self.check_grad_with_place( - place, - {'X'}, - 'Out', - check_dygraph=(not self.use_mkldnn), - check_cinn=True, - check_pir=True, - check_pir_onednn=self.check_pir_onednn, - ) - elif self.pool_type != "max": - self.check_grad( - {'X'}, - 'Out', - max_relative_error=0.07, - check_dygraph=(not self.use_mkldnn), - check_pir=True, - check_pir_onednn=self.check_pir_onednn, - ) - - def init_data_format(self): - self.data_format = "NCHW" - - def init_shape(self): - self.shape = [2, 3, 5, 5] - - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - - def init_paddings(self): - self.paddings = [0, 0] - self.padding_algorithm = "EXPLICIT" - - def init_kernel_type(self): - self.use_cudnn = False - - def init_data_type(self): - self.dtype = np.float32 if core.is_compiled_with_rocm() else np.float64 - - def init_pool_type(self): - self.pool_type = "avg" - self.pool2D_forward_naive = avg_pool2D_forward_naive - - def init_global_pool(self): - self.global_pool = True - - def init_ceil_mode(self): - self.ceil_mode = False - - def init_exclusive(self): - self.exclusive = True - - def init_adaptive(self): - self.adaptive = False - - -class TestPool2D_Op(TestPool2D_Op_Mixin, OpTest): - pass - - -class TestCase1(TestPool2D_Op): - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - - def init_paddings(self): - self.paddings = [0, 0] - - def init_pool_type(self): - self.pool_type = "avg" - self.pool2D_forward_naive = avg_pool2D_forward_naive - - def init_global_pool(self): - self.global_pool = False - - def init_shape(self): - self.shape = [2, 3, 7, 7] - - -class TestCase2(TestPool2D_Op): - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - - def init_paddings(self): - self.paddings = [1, 1] - - def init_pool_type(self): - self.pool_type = "avg" - self.pool2D_forward_naive = avg_pool2D_forward_naive - - def init_global_pool(self): - self.global_pool = False - - def init_shape(self): - self.shape = [2, 3, 7, 7] - - -class TestCase3(TestPool2D_Op): - def init_pool_type(self): - self.pool_type = "max" - self.pool2D_forward_naive = max_pool2D_forward_naive - - -class TestCase4(TestCase1): - def init_pool_type(self): - self.pool_type = "max" - self.pool2D_forward_naive = max_pool2D_forward_naive - - -class TestCase5(TestCase2): - def init_pool_type(self): - self.pool_type = "max" - self.pool2D_forward_naive = max_pool2D_forward_naive - - -# --------------------test pool2d cudnn-------------------- - - -def create_test_cudnn_class(parent): - @unittest.skipIf( - not core.is_compiled_with_cuda(), "core is not compiled with CUDA" - ) - class TestCUDNNCase(parent): - def init_kernel_type(self): - self.use_cudnn = True - - cls_name = "{}_{}".format(parent.__name__, "CUDNNOp") - TestCUDNNCase.__name__ = cls_name - globals()[cls_name] = TestCUDNNCase - - -create_test_cudnn_class(TestPool2D_Op) -create_test_cudnn_class(TestCase1) -create_test_cudnn_class(TestCase2) -create_test_cudnn_class(TestCase3) -create_test_cudnn_class(TestCase4) -create_test_cudnn_class(TestCase5) - -# --------------------test pool2d cudnn_fp16-------------------- - - -def create_test_cudnn_fp16_class(parent, check_grad=True): - @unittest.skipIf( - not core.is_compiled_with_cuda(), "core is not compiled with CUDA" - ) - class TestCUDNNFp16Case(parent): - def init_kernel_type(self): - self.use_cudnn = True - self.dtype = np.float16 - - def test_check_output(self): - # TODO(wangzhongpu): support onednn op in dygraph mode - if core.is_compiled_with_cuda(): - place = core.CUDAPlace(0) - if core.is_float16_supported(place): - self.check_output_with_place( - place, - check_dygraph=(not self.use_mkldnn), - check_cinn=True, - check_pir_onednn=self.check_pir_onednn, - ) - - def test_check_grad(self): - # TODO(wangzhongpu): support onednn op in dygraph mode - place = core.CUDAPlace(0) - if ( - core.is_float16_supported(place) - and self.pool_type != "max" - and check_grad - ): - self.check_grad_with_place( - place, - {'X'}, - 'Out', - check_dygraph=(not self.use_mkldnn), - check_cinn=True, - check_pir_onednn=self.check_pir_onednn, - ) - - cls_name = "{}_{}".format(parent.__name__, "CUDNNFp16Op") - TestCUDNNFp16Case.__name__ = cls_name - globals()[cls_name] = TestCUDNNFp16Case - - -def create_test_fp16_class(parent, check_grad=True): - @unittest.skipIf( - not core.is_compiled_with_cuda(), "core is not compiled with CUDA" - ) - class TestFp16Case(parent): - def init_kernel_type(self): - self.use_cudnn = False - self.dtype = np.float16 - - def test_check_output(self): - # TODO(wangzhongpu): support onednn op in dygraph mode - if core.is_compiled_with_cuda(): - place = core.CUDAPlace(0) - if core.is_float16_supported(place): - self.check_output_with_place( - place, - check_dygraph=(not self.use_mkldnn), - check_cinn=True, - check_pir_onednn=self.check_pir_onednn, - ) - - def test_check_grad(self): - # TODO(wangzhongpu): support onednn op in dygraph mode - place = core.CUDAPlace(0) - if ( - core.is_float16_supported(place) - and self.pool_type != "max" - and check_grad - ): - self.check_grad_with_place( - place, - {'X'}, - 'Out', - check_dygraph=(not self.use_mkldnn), - check_cinn=True, - check_pir_onednn=self.check_pir_onednn, - ) - - cls_name = "{}_{}".format(parent.__name__, "Fp16Op") - TestFp16Case.__name__ = cls_name - globals()[cls_name] = TestFp16Case - - -def create_test_bf16_class(parent, check_grad=True): - @unittest.skipIf( - not core.is_compiled_with_cuda(), "core is not compiled with CUDA" - ) - class TestBf16Case(parent): - def init_kernel_type(self): - self.use_cuda = True - self.dtype = np.uint16 - - def test_check_output(self): - if core.is_compiled_with_cuda(): - place = core.CUDAPlace(0) - self.check_output_with_place( - place, - check_dygraph=(not self.use_mkldnn), - check_cinn=True, - check_pir_onednn=self.check_pir_onednn, - ) - - def test_check_grad(self): - place = core.CUDAPlace(0) - if self.pool_type != "max" and check_grad: - self.check_grad_with_place( - place, - {'X'}, - 'Out', - check_dygraph=(not self.use_mkldnn), - check_cinn=True, - check_pir_onednn=self.check_pir_onednn, - ) - - cls_name = "{}_{}".format(parent.__name__, "Bf16Op") - TestBf16Case.__name__ = cls_name - globals()[cls_name] = TestBf16Case - - -create_test_cudnn_fp16_class(TestPool2D_Op) -create_test_cudnn_fp16_class(TestCase1) -create_test_cudnn_fp16_class(TestCase2) -create_test_cudnn_fp16_class(TestCase3) -create_test_cudnn_fp16_class(TestCase4) -create_test_cudnn_fp16_class(TestCase5) - -create_test_fp16_class(TestPool2D_Op) -create_test_fp16_class(TestCase1) -create_test_fp16_class(TestCase2) -create_test_fp16_class(TestCase3) -create_test_fp16_class(TestCase4) -create_test_fp16_class(TestCase5) - -create_test_bf16_class(TestPool2D_Op) -create_test_bf16_class(TestCase1) -create_test_bf16_class(TestCase2) -create_test_bf16_class(TestCase3) -create_test_bf16_class(TestCase4) -create_test_bf16_class(TestCase5) -# --------------------test pool2d use ceil mode-------------------- - - -def create_test_cudnn_use_ceil_class(parent): - @unittest.skipIf( - not core.is_compiled_with_cuda(), "core is not compiled with CUDA" - ) - class TestPool2DUseCeilCase(parent): - def init_kernel_type(self): - self.use_cudnn = True - - def init_ceil_mode(self): - self.ceil_mode = True - - cls_name = "{}_{}".format(parent.__name__, "CUDNNOpCeilMode") - TestPool2DUseCeilCase.__name__ = cls_name - globals()[cls_name] = TestPool2DUseCeilCase - - -create_test_cudnn_use_ceil_class(TestPool2D_Op) -create_test_cudnn_use_ceil_class(TestCase1) - - -def create_test_use_ceil_class(parent): - class TestPool2DUseCeilCase(parent): - def init_ceil_mode(self): - self.ceil_mode = True - - cls_name = "{}_{}".format(parent.__name__, "CeilModeCast") - TestPool2DUseCeilCase.__name__ = cls_name - globals()[cls_name] = TestPool2DUseCeilCase - - -create_test_use_ceil_class(TestCase1) -create_test_use_ceil_class(TestCase2) - - -class TestAvgInclude(TestCase2): - def init_exclusive(self): - self.exclusive = False - - -class TestCUDNNAvgInclude(TestCase2): - def init_kernel_type(self): - self.use_cudnn = True - - def init_exclusive(self): - self.exclusive = False - - -class TestAvgPoolAdaptive(TestCase1): - def init_adaptive(self): - self.adaptive = True - - -class TestAvgPoolAdaptiveAsyOutSize(TestCase1): - def init_adaptive(self): - self.adaptive = True - - def init_shape(self): - self.shape = [8, 3, 6, 6] - - def init_test_case(self): - self.ksize = [2, 3] - self.strides = [1, 1] - self.paddings = [0, 0, 0, 0] - - -# -------test pool2d with asymmetric padding----- - - -class TestPool2D_AsyPadding(TestPool2D_Op): - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - self.paddings = [1, 0, 1, 2] - - def init_shape(self): - self.shape = [2, 3, 5, 5] - - -class TestCase1_AsyPadding(TestCase1): - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - self.paddings = [1, 0, 1, 0] - - def init_shape(self): - self.shape = [2, 3, 7, 7] - - -class TestCase2_AsyPadding(TestCase2): - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - self.paddings = [1, 2, 1, 2] - - def init_shape(self): - self.shape = [2, 3, 7, 7] - - -class TestCase3_AsyPadding(TestCase3): - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - self.paddings = [1, 0, 1, 2] - - def init_shape(self): - self.shape = [2, 3, 5, 5] - - -class TestCase4_AsyPadding(TestCase4): - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - self.paddings = [1, 0, 1, 0] - - def init_shape(self): - self.shape = [2, 3, 7, 7] - - -class TestCase5_AsyPadding(TestCase5): - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - self.paddings = [2, 2, 1, 2] - - def init_shape(self): - self.shape = [2, 3, 7, 7] - - -create_test_cudnn_class(TestPool2D_AsyPadding) -create_test_cudnn_class(TestCase1_AsyPadding) -create_test_cudnn_class(TestCase2_AsyPadding) -create_test_cudnn_class(TestCase3_AsyPadding) -create_test_cudnn_class(TestCase4_AsyPadding) -create_test_cudnn_class(TestCase5_AsyPadding) - -create_test_cudnn_fp16_class(TestPool2D_AsyPadding) -create_test_cudnn_fp16_class(TestCase1_AsyPadding) -create_test_cudnn_fp16_class(TestCase2_AsyPadding) -create_test_cudnn_fp16_class(TestCase3_AsyPadding) -create_test_cudnn_fp16_class(TestCase4_AsyPadding) -create_test_cudnn_fp16_class(TestCase5_AsyPadding) - -create_test_fp16_class(TestPool2D_AsyPadding) -create_test_fp16_class(TestCase1_AsyPadding) -create_test_fp16_class(TestCase2_AsyPadding) -create_test_fp16_class(TestCase3_AsyPadding) -create_test_fp16_class(TestCase4_AsyPadding) -create_test_fp16_class(TestCase5_AsyPadding) - -create_test_bf16_class(TestPool2D_AsyPadding) -create_test_bf16_class(TestCase1_AsyPadding) -create_test_bf16_class(TestCase2_AsyPadding) -create_test_bf16_class(TestCase3_AsyPadding) -create_test_bf16_class(TestCase4_AsyPadding) -create_test_bf16_class(TestCase5_AsyPadding) - -create_test_cudnn_use_ceil_class(TestPool2D_AsyPadding) -create_test_cudnn_use_ceil_class(TestCase1_AsyPadding) - -create_test_use_ceil_class(TestCase1_AsyPadding) -create_test_use_ceil_class(TestCase2_AsyPadding) - - -class TestAvgInclude_AsyPadding(TestCase2): - def init_exclusive(self): - self.exclusive = False - - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - self.paddings = [1, 2, 1, 2] - - def init_shape(self): - self.shape = [2, 3, 7, 7] - - -class TestCUDNNAvgInclude_AsyPadding(TestCase2): - def init_kernel_type(self): - self.use_cudnn = True - - def init_exclusive(self): - self.exclusive = False - - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - self.paddings = [2, 1, 1, 1] - - def init_shape(self): - self.shape = [2, 3, 7, 7] - - -class TestAvgPoolAdaptive_AsyPadding(TestCase1): - def init_adaptive(self): - self.adaptive = True - - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 1] - self.paddings = [1, 1, 0, 2] - - def init_shape(self): - self.shape = [2, 3, 7, 7] - - -# ----------- test channel_last -------------- -class TestPool2D_channel_last(TestPool2D_Op): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 5, 5, 3] - - -class TestCase1_channel_last(TestCase1): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -class TestCase2_channel_last(TestCase2): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -class TestCase3_channel_last(TestCase3): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 5, 5, 3] - - -class TestCase4_channel_last(TestCase4): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -class TestCase5_channel_last(TestCase5): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -create_test_cudnn_class(TestPool2D_channel_last) -create_test_cudnn_class(TestCase1_channel_last) -create_test_cudnn_class(TestCase2_channel_last) -create_test_cudnn_class(TestCase3_channel_last) -create_test_cudnn_class(TestCase4_channel_last) -create_test_cudnn_class(TestCase5_channel_last) - -create_test_cudnn_fp16_class(TestPool2D_channel_last) -create_test_cudnn_fp16_class(TestCase1_channel_last) -create_test_cudnn_fp16_class(TestCase2_channel_last) -create_test_cudnn_fp16_class(TestCase3_channel_last) -create_test_cudnn_fp16_class(TestCase4_channel_last) -create_test_cudnn_fp16_class(TestCase5_channel_last) - -create_test_fp16_class(TestPool2D_channel_last) -create_test_fp16_class(TestCase1_channel_last) -create_test_fp16_class(TestCase2_channel_last) -create_test_fp16_class(TestCase3_channel_last) -create_test_fp16_class(TestCase4_channel_last) -create_test_fp16_class(TestCase5_channel_last) - -create_test_bf16_class(TestPool2D_channel_last) -create_test_bf16_class(TestCase1_channel_last) -create_test_bf16_class(TestCase2_channel_last) -create_test_bf16_class(TestCase3_channel_last) -create_test_bf16_class(TestCase4_channel_last) -create_test_bf16_class(TestCase5_channel_last) - -create_test_cudnn_use_ceil_class(TestPool2D_channel_last) -create_test_cudnn_use_ceil_class(TestCase1_channel_last) - -create_test_use_ceil_class(TestCase1_channel_last) -create_test_use_ceil_class(TestCase2_channel_last) - - -class TestCase5_Max(TestCase2): - def init_pool_type(self): - self.pool_type = "max" - - def test_check_grad(self): - if self.dtype == np.float16: - return - if self.has_cudnn() and self.pool_type == "max": - place = core.CUDAPlace(0) - self.check_grad_with_place( - place, - {'X'}, - 'Out', - max_relative_error=1.00, - check_cinn=True, - check_pir=True, - check_pir_onednn=self.check_pir_onednn, - ) - elif self.pool_type == "max": - self.check_grad( - {'X'}, - 'Out', - max_relative_error=1.00, - check_cinn=True, - check_pir=True, - check_pir_onednn=self.check_pir_onednn, - ) - - -class TestCase5_channel_last_Max(TestCase5_Max): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -create_test_cudnn_class(TestCase5_Max) -create_test_cudnn_class(TestCase5_channel_last_Max) - - -class TestAvgInclude_channel_last(TestCase2_channel_last): - def init_exclusive(self): - self.exclusive = False - - -class TestCUDNNAvgInclude_channel_last(TestCase2_channel_last): - def init_kernel_type(self): - self.use_cudnn = True - - def init_exclusive(self): - self.exclusive = False - - -class TestAvgPoolAdaptive_channel_last(TestCase1_channel_last): - def init_adaptive(self): - self.adaptive = True - - -class TestPool2D_AsyPadding_channel_last(TestPool2D_AsyPadding): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 5, 5, 3] - - -class TestCase1_AsyPadding_channel_last(TestCase1_AsyPadding): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -class TestCase2_AsyPadding_channel_last(TestCase2_AsyPadding): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -class TestCase3_AsyPadding_channel_last(TestCase3_AsyPadding): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 5, 5, 3] - - -class TestCase4_AsyPadding_channel_last(TestCase4_AsyPadding): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -class TestCase5_AsyPadding_channel_last(TestCase5_AsyPadding): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -create_test_cudnn_class(TestPool2D_AsyPadding_channel_last) -create_test_cudnn_class(TestCase1_AsyPadding_channel_last) -create_test_cudnn_class(TestCase2_AsyPadding_channel_last) -create_test_cudnn_class(TestCase3_AsyPadding_channel_last) -create_test_cudnn_class(TestCase4_AsyPadding_channel_last) -create_test_cudnn_class(TestCase5_AsyPadding_channel_last) - -create_test_cudnn_fp16_class(TestPool2D_AsyPadding_channel_last) -create_test_cudnn_fp16_class(TestCase1_AsyPadding_channel_last) -create_test_cudnn_fp16_class(TestCase2_AsyPadding_channel_last) -create_test_cudnn_fp16_class(TestCase3_AsyPadding_channel_last) -create_test_cudnn_fp16_class(TestCase4_AsyPadding_channel_last) -create_test_cudnn_fp16_class(TestCase5_AsyPadding_channel_last) - -create_test_fp16_class(TestPool2D_AsyPadding_channel_last) -create_test_fp16_class(TestCase1_AsyPadding_channel_last) -create_test_fp16_class(TestCase2_AsyPadding_channel_last) -create_test_fp16_class(TestCase3_AsyPadding_channel_last) -create_test_fp16_class(TestCase4_AsyPadding_channel_last) -create_test_fp16_class(TestCase5_AsyPadding_channel_last) - -create_test_bf16_class(TestPool2D_AsyPadding_channel_last) -create_test_bf16_class(TestCase1_AsyPadding_channel_last) -create_test_bf16_class(TestCase2_AsyPadding_channel_last) -create_test_bf16_class(TestCase3_AsyPadding_channel_last) -create_test_bf16_class(TestCase4_AsyPadding_channel_last) -create_test_bf16_class(TestCase5_AsyPadding_channel_last) - -create_test_cudnn_use_ceil_class(TestPool2D_AsyPadding_channel_last) -create_test_cudnn_use_ceil_class(TestCase1_AsyPadding_channel_last) - -create_test_use_ceil_class(TestCase1_AsyPadding_channel_last) -create_test_use_ceil_class(TestCase2_AsyPadding_channel_last) - - -class TestAvgInclude_AsyPadding_channel_last(TestAvgInclude_AsyPadding): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -class TestCUDNNAvgInclude_AsyPadding_channel_last( - TestCUDNNAvgInclude_AsyPadding -): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -class TestAvgPoolAdaptive_AsyPadding_channel_last( - TestAvgPoolAdaptive_AsyPadding -): - def init_data_format(self): - self.data_format = "NHWC" - - def init_shape(self): - self.shape = [2, 7, 7, 3] - - -# test paddings: SAME VALID - - -def create_test_padding_SAME_class(parent): - class TestPaddingSMAECase(parent): - def init_paddings(self): - self.paddings = [0, 0] - self.padding_algorithm = "SAME" - - cls_name = "{}_{}".format(parent.__name__, "PaddingSAMEOp") - TestPaddingSMAECase.__name__ = cls_name - globals()[cls_name] = TestPaddingSMAECase - - -create_test_padding_SAME_class(TestPool2D_Op) -create_test_padding_SAME_class(TestCase1) -create_test_padding_SAME_class(TestCase2) -create_test_padding_SAME_class(TestCase3) -create_test_padding_SAME_class(TestCase4) -create_test_padding_SAME_class(TestCase5) - -create_test_padding_SAME_class(TestPool2D_channel_last) -create_test_padding_SAME_class(TestCase1_channel_last) -create_test_padding_SAME_class(TestCase2_channel_last) -create_test_padding_SAME_class(TestCase3_channel_last) -create_test_padding_SAME_class(TestCase4_channel_last) -create_test_padding_SAME_class(TestCase5_channel_last) - - -def create_test_cudnn_padding_SAME_class(parent): - @unittest.skipIf( - not core.is_compiled_with_cuda(), "core is not compiled with CUDA" - ) - class TestCUDNNPaddingSMAECase(parent): - def init_kernel_type(self): - self.use_cudnn = True - - def init_paddings(self): - self.paddings = [1, 1] - self.padding_algorithm = "SAME" - - cls_name = "{}_{}".format(parent.__name__, "CudnnPaddingSAMEOp") - TestCUDNNPaddingSMAECase.__name__ = cls_name - globals()[cls_name] = TestCUDNNPaddingSMAECase - - -create_test_cudnn_padding_SAME_class(TestPool2D_Op) -create_test_cudnn_padding_SAME_class(TestCase1) -create_test_cudnn_padding_SAME_class(TestCase2) -create_test_cudnn_padding_SAME_class(TestCase3) -create_test_cudnn_padding_SAME_class(TestCase4) -create_test_cudnn_padding_SAME_class(TestCase5) - -create_test_cudnn_padding_SAME_class(TestPool2D_channel_last) -create_test_cudnn_padding_SAME_class(TestCase1_channel_last) -create_test_cudnn_padding_SAME_class(TestCase2_channel_last) -create_test_cudnn_padding_SAME_class(TestCase3_channel_last) -create_test_cudnn_padding_SAME_class(TestCase4_channel_last) -create_test_cudnn_padding_SAME_class(TestCase5_channel_last) - - -def create_test_padding_VALID_class(parent): - class TestPaddingVALIDCase(parent): - def init_paddings(self): - self.paddings = [1, 1] - self.padding_algorithm = "VALID" - - cls_name = "{}_{}".format(parent.__name__, "PaddingVALIDOp") - TestPaddingVALIDCase.__name__ = cls_name - globals()[cls_name] = TestPaddingVALIDCase - - -create_test_padding_VALID_class(TestPool2D_Op) -create_test_padding_VALID_class(TestCase1) -create_test_padding_VALID_class(TestCase2) -create_test_padding_VALID_class(TestCase3) -create_test_padding_VALID_class(TestCase4) -create_test_padding_VALID_class(TestCase5) - -create_test_padding_VALID_class(TestPool2D_channel_last) -create_test_padding_VALID_class(TestCase1_channel_last) -create_test_padding_VALID_class(TestCase2_channel_last) -create_test_padding_VALID_class(TestCase3_channel_last) -create_test_padding_VALID_class(TestCase4_channel_last) -create_test_padding_VALID_class(TestCase5_channel_last) - - -def create_test_cudnn_padding_VALID_class(parent): - @unittest.skipIf( - not core.is_compiled_with_cuda(), "core is not compiled with CUDA" - ) - class TestCUDNNPaddingVALIDCase(parent): - def init_kernel_type(self): - self.use_cudnn = True - - def init_paddings(self): - self.paddings = [1, 1] - self.padding_algorithm = "VALID" - - cls_name = "{}_{}".format(parent.__name__, "CudnnPaddingVALIDOp") - TestCUDNNPaddingVALIDCase.__name__ = cls_name - globals()[cls_name] = TestCUDNNPaddingVALIDCase - - -create_test_cudnn_padding_VALID_class(TestPool2D_Op) -create_test_cudnn_padding_VALID_class(TestCase1) -create_test_cudnn_padding_VALID_class(TestCase2) -create_test_cudnn_padding_VALID_class(TestCase3) -create_test_cudnn_padding_VALID_class(TestCase4) -create_test_cudnn_padding_VALID_class(TestCase5) - -create_test_cudnn_padding_VALID_class(TestPool2D_channel_last) -create_test_cudnn_padding_VALID_class(TestCase1_channel_last) -create_test_cudnn_padding_VALID_class(TestCase2_channel_last) -create_test_cudnn_padding_VALID_class(TestCase3_channel_last) -create_test_cudnn_padding_VALID_class(TestCase4_channel_last) -create_test_cudnn_padding_VALID_class(TestCase5_channel_last) - - -class TestCase1_strides(TestCase1): - def init_test_case(self): - self.ksize = [3, 3] - self.strides = [1, 2] - - def init_shape(self): - self.shape = [2, 3, 4, 5] - - -create_test_cudnn_class(TestCase1_strides) -create_test_padding_SAME_class(TestCase1_strides) -create_test_cudnn_padding_SAME_class(TestCase1_strides) - - -if __name__ == '__main__': - unittest.main() From 6fcf989efcbefe3cd317831ea5286930d211e9ba Mon Sep 17 00:00:00 2001 From: Wang Huan Date: Tue, 14 May 2024 03:19:01 +0000 Subject: [PATCH 2/6] refine --- ...est_elementwise_gradient_op_deprecated.py} | 0 ...test_inplace_addto_strategy_deprecated.py} | 0 ...tr.py => test_set_bool_attr_deprecated.py} | 1 + .../test_fractional_max_pool3d_op.py | 0 test/legacy_test/test_pool2d_op.py | 1326 +++++++++++++++++ 5 files changed, 1327 insertions(+) rename test/deprecated/legacy_test/{test_elementwise_gradient_op.py => test_elementwise_gradient_op_deprecated.py} (100%) rename test/deprecated/legacy_test/{test_inplace_addto_strategy.py => test_inplace_addto_strategy_deprecated.py} (100%) rename test/deprecated/legacy_test/{test_set_bool_attr.py => test_set_bool_attr_deprecated.py} (98%) rename test/{deprecated => }/legacy_test/test_fractional_max_pool3d_op.py (100%) create mode 100644 test/legacy_test/test_pool2d_op.py diff --git a/test/deprecated/legacy_test/test_elementwise_gradient_op.py b/test/deprecated/legacy_test/test_elementwise_gradient_op_deprecated.py similarity index 100% rename from test/deprecated/legacy_test/test_elementwise_gradient_op.py rename to test/deprecated/legacy_test/test_elementwise_gradient_op_deprecated.py diff --git a/test/deprecated/legacy_test/test_inplace_addto_strategy.py b/test/deprecated/legacy_test/test_inplace_addto_strategy_deprecated.py similarity index 100% rename from test/deprecated/legacy_test/test_inplace_addto_strategy.py rename to test/deprecated/legacy_test/test_inplace_addto_strategy_deprecated.py diff --git a/test/deprecated/legacy_test/test_set_bool_attr.py b/test/deprecated/legacy_test/test_set_bool_attr_deprecated.py similarity index 98% rename from test/deprecated/legacy_test/test_set_bool_attr.py rename to test/deprecated/legacy_test/test_set_bool_attr_deprecated.py index 5f7282e12a1f89..1752a68f09e089 100644 --- a/test/deprecated/legacy_test/test_set_bool_attr.py +++ b/test/deprecated/legacy_test/test_set_bool_attr_deprecated.py @@ -44,4 +44,5 @@ def test_set_bool_attr(self): if __name__ == '__main__': + paddle.enable_static() unittest.main() diff --git a/test/deprecated/legacy_test/test_fractional_max_pool3d_op.py b/test/legacy_test/test_fractional_max_pool3d_op.py similarity index 100% rename from test/deprecated/legacy_test/test_fractional_max_pool3d_op.py rename to test/legacy_test/test_fractional_max_pool3d_op.py diff --git a/test/legacy_test/test_pool2d_op.py b/test/legacy_test/test_pool2d_op.py new file mode 100644 index 00000000000000..b2f10e3af1b266 --- /dev/null +++ b/test/legacy_test/test_pool2d_op.py @@ -0,0 +1,1326 @@ +# Copyright (c) 2018 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. + +import unittest + +import numpy as np +from op_test import OpTest, convert_float_to_uint16 + +import paddle +from paddle.base import core +from paddle.framework import in_dynamic_mode + + +def adaptive_start_index(index, input_size, output_size): + return int(np.floor(index * input_size / output_size)) + + +def adaptive_end_index(index, input_size, output_size): + return int(np.ceil((index + 1) * input_size / output_size)) + + +def max_pool2D_forward_naive( + x, + ksize, + strides, + paddings, + global_pool=0, + ceil_mode=False, + exclusive=True, + adaptive=False, + data_type=np.float64, +): + if data_type == np.float64 and core.is_compiled_with_rocm(): + data_type = np.float32 + N, C, H, W = x.shape + if global_pool == 1: + ksize = [H, W] + if adaptive: + H_out, W_out = ksize + else: + H_out = ( + (H - ksize[0] + 2 * paddings[0] + strides[0] - 1) // strides[0] + 1 + if ceil_mode + else (H - ksize[0] + 2 * paddings[0]) // strides[0] + 1 + ) + W_out = ( + (W - ksize[1] + 2 * paddings[1] + strides[1] - 1) // strides[1] + 1 + if ceil_mode + else (W - ksize[1] + 2 * paddings[1]) // strides[1] + 1 + ) + out = np.zeros((N, C, H_out, W_out)) + for i in range(H_out): + for j in range(W_out): + if adaptive: + r_start = adaptive_start_index(i, H, ksize[0]) + r_end = adaptive_end_index(i, H, ksize[0]) + c_start = adaptive_start_index(j, W, ksize[1]) + c_end = adaptive_end_index(j, W, ksize[1]) + else: + r_start = np.max((i * strides[0] - paddings[0], 0)) + r_end = np.min((i * strides[0] + ksize[0] - paddings[0], H)) + c_start = np.max((j * strides[1] - paddings[1], 0)) + c_end = np.min((j * strides[1] + ksize[1] - paddings[1], W)) + x_masked = x[:, :, r_start:r_end, c_start:c_end] + + out[:, :, i, j] = np.max(x_masked, axis=(2, 3)) + return out + + +def avg_pool2D_forward_naive( + x, + ksize, + strides, + paddings, + global_pool=0, + ceil_mode=False, + exclusive=True, + adaptive=False, + data_type=np.float64, +): + if data_type == np.float64 and core.is_compiled_with_rocm(): + data_type = np.float32 + N, C, H, W = x.shape + if global_pool == 1: + ksize = [H, W] + if adaptive: + H_out, W_out = ksize + else: + H_out = ( + (H - ksize[0] + 2 * paddings[0] + strides[0] - 1) // strides[0] + 1 + if ceil_mode + else (H - ksize[0] + 2 * paddings[0]) // strides[0] + 1 + ) + W_out = ( + (W - ksize[1] + 2 * paddings[1] + strides[1] - 1) // strides[1] + 1 + if ceil_mode + else (W - ksize[1] + 2 * paddings[1]) // strides[1] + 1 + ) + out = np.zeros((N, C, H_out, W_out)) + for i in range(H_out): + for j in range(W_out): + if adaptive: + r_start = adaptive_start_index(i, H, ksize[0]) + r_end = adaptive_end_index(i, H, ksize[0]) + c_start = adaptive_start_index(j, W, ksize[1]) + c_end = adaptive_end_index(j, W, ksize[1]) + else: + r_start = i * strides[0] - paddings[0] + r_end = i * strides[0] + ksize[0] - paddings[0] + c_start = j * strides[1] - paddings[1] + c_end = j * strides[1] + ksize[1] - paddings[1] + field_size = (r_end - r_start) * (c_end - c_start) + r_start = np.max((r_start, 0)) + r_end = np.min((r_end, H)) + c_start = np.max((c_start, 0)) + c_end = np.min((c_end, W)) + + x_masked = x[:, :, r_start:r_end, c_start:c_end] + + if exclusive or adaptive: + field_size = (r_end - r_start) * (c_end - c_start) + + if data_type == np.int8 or data_type == np.uint8: + out[:, :, i, j] = ( + np.rint(np.sum(x_masked, axis=(2, 3)) / field_size) + ).astype(data_type) + else: + out[:, :, i, j] = ( + np.sum(x_masked, axis=(2, 3)) / field_size + ).astype(data_type) + return out + + +def pool2D_forward_naive( + x, + ksize, + strides, + paddings, + global_pool=0, + ceil_mode=False, + exclusive=True, + adaptive=False, + data_format='NCHW', + pool_type="max", + padding_algorithm="EXPLICIT", +): + # update paddings + def _get_padding_with_SAME(input_shape, pool_size, pool_stride): + padding = [] + for input_size, filter_size, stride_size in zip( + input_shape, pool_size, pool_stride + ): + out_size = int((input_size + stride_size - 1) / stride_size) + pad_sum = np.max( + ((out_size - 1) * stride_size + filter_size - input_size, 0) + ) + pad_0 = int(pad_sum / 2) + pad_1 = int(pad_sum - pad_0) + padding.append(pad_0) + padding.append(pad_1) + return padding + + if isinstance(padding_algorithm, str): + padding_algorithm = padding_algorithm.upper() + if padding_algorithm not in ["SAME", "VALID", "EXPLICIT"]: + raise ValueError( + "Unknown Attr(padding_algorithm): '%s'. " + "It can only be 'SAME' or 'VALID'." % str(padding_algorithm) + ) + + if padding_algorithm == "VALID": + paddings = [0, 0, 0, 0] + if ceil_mode is not False: + raise ValueError( + "When Attr(pool_padding) is \"VALID\", Attr(ceil_mode)" + " must be False. " + "Received ceil_mode: True." + ) + elif padding_algorithm == "SAME": + input_data_shape = [] + if data_format == "NCHW": + input_data_shape = x.shape[2:4] + elif data_format == "NHWC": + input_data_shape = x.shape[1:3] + paddings = _get_padding_with_SAME(input_data_shape, ksize, strides) + + assert len(paddings) == 2 or len(paddings) == 4 + is_sys = True if len(paddings) == 2 else False + + N = x.shape[0] + C, H, W = ( + [x.shape[1], x.shape[2], x.shape[3]] + if data_format == 'NCHW' + else [x.shape[3], x.shape[1], x.shape[2]] + ) + + if global_pool == 1: + ksize = [H, W] + paddings = [0 for _ in range(len(paddings))] + + pad_h_up = paddings[0] if is_sys else paddings[0] + pad_h_down = paddings[0] if is_sys else paddings[1] + pad_w_left = paddings[1] if is_sys else paddings[2] + pad_w_right = paddings[1] if is_sys else paddings[3] + + if adaptive: + H_out, W_out = ksize + else: + H_out = ( + (H - ksize[0] + pad_h_up + pad_h_down + strides[0] - 1) + // strides[0] + + 1 + if ceil_mode + else (H - ksize[0] + pad_h_up + pad_h_down) // strides[0] + 1 + ) + W_out = ( + (W - ksize[1] + pad_w_left + pad_w_right + strides[1] - 1) + // strides[1] + + 1 + if ceil_mode + else (W - ksize[1] + pad_w_left + pad_w_right) // strides[1] + 1 + ) + + out = ( + np.zeros((N, C, H_out, W_out)) + if data_format == 'NCHW' + else np.zeros((N, H_out, W_out, C)) + ) + for i in range(H_out): + if adaptive: + in_h_start = adaptive_start_index(i, H, ksize[0]) + in_h_end = adaptive_end_index(i, H, ksize[0]) + else: + in_h_start = np.max((i * strides[0] - pad_h_up, 0)) + in_h_end = np.min((i * strides[0] + ksize[0] - pad_h_up, H)) + + for j in range(W_out): + if adaptive: + in_w_start = adaptive_start_index(j, W, ksize[1]) + in_w_end = adaptive_end_index(j, W, ksize[1]) + else: + in_h_start = i * strides[0] - pad_h_up + in_w_start = j * strides[1] - pad_w_left + in_h_end = i * strides[0] + ksize[0] - pad_h_up + in_w_end = j * strides[1] + ksize[1] - pad_w_left + + field_size = (in_h_end - in_h_start) * (in_w_end - in_w_start) + in_h_start = np.max((in_h_start, 0)) + in_w_start = np.max((in_w_start, 0)) + in_h_end = np.min((in_h_end, H)) + in_w_end = np.min((in_w_end, W)) + + if data_format == 'NCHW': + x_masked = x[:, :, in_h_start:in_h_end, in_w_start:in_w_end] + if pool_type == 'avg': + if exclusive or adaptive: + field_size = (in_h_end - in_h_start) * ( + in_w_end - in_w_start + ) + + # if (exclusive or adaptive) else (ksize[0] * ksize[1]) + out[:, :, i, j] = np.sum(x_masked, axis=(2, 3)) / field_size + elif pool_type == 'max': + out[:, :, i, j] = np.max(x_masked, axis=(2, 3)) + elif data_format == 'NHWC': + x_masked = x[:, in_h_start:in_h_end, in_w_start:in_w_end, :] + if pool_type == 'avg': + if exclusive or adaptive: + field_size = (in_h_end - in_h_start) * ( + in_w_end - in_w_start + ) + out[:, i, j, :] = np.sum(x_masked, axis=(1, 2)) / field_size + elif pool_type == 'max': + out[:, i, j, :] = np.max(x_masked, axis=(1, 2)) + return out + + +def pool2d_wrapper_not_use_cudnn( + X, + ksize=[], + strides=[], + paddings=[], + ceil_mode=False, + exclusive=True, + data_format="NCDHW", + pooling_type="max", + global_pooling=False, + adaptive=False, + padding_algorithm="EXPLICIT", +): + if in_dynamic_mode(): + X = X._use_gpudnn(False) + if data_format == "AnyLayout": + data_format = "NCDHW" + return paddle._C_ops.pool2d( + X, + ksize, + strides, + paddings, + ceil_mode, + exclusive, + data_format, + pooling_type, + global_pooling, + adaptive, + padding_algorithm, + ) + + +def pool2d_wrapper_use_cudnn( + X, + ksize=[], + strides=[], + paddings=[], + ceil_mode=False, + exclusive=True, + data_format="NCDHW", + pooling_type="max", + global_pooling=False, + adaptive=False, + padding_algorithm="EXPLICIT", +): + if data_format == "AnyLayout": + data_format = "NCDHW" + return paddle._C_ops.pool2d( + X, + ksize, + strides, + paddings, + ceil_mode, + exclusive, + data_format, + pooling_type, + global_pooling, + adaptive, + padding_algorithm, + ) + + +class TestPool2D_Op_Mixin: + def setUp(self): + self.op_type = "pool2d" + self.use_cudnn = False + self.init_kernel_type() + self.use_mkldnn = False + self.init_data_type() + self.init_test_case() + self.padding_algorithm = "EXPLICIT" + self.init_paddings() + self.init_global_pool() + self.init_kernel_type() + self.init_pool_type() + self.init_ceil_mode() + self.init_exclusive() + self.init_adaptive() + self.init_data_format() + self.init_shape() + + if self.is_bfloat16_op(): + input = np.random.random(self.shape).astype(np.float32) + else: + input = np.random.random(self.shape).astype(self.dtype) + + output = pool2D_forward_naive( + input, + self.ksize, + self.strides, + self.paddings, + self.global_pool, + self.ceil_mode, + self.exclusive, + self.adaptive, + self.data_format, + self.pool_type, + self.padding_algorithm, + ) + + if self.is_bfloat16_op(): + output = convert_float_to_uint16(output) + self.inputs = {'X': convert_float_to_uint16(input)} + else: + output = output.astype(self.dtype) + self.inputs = {'X': OpTest.np_dtype_to_base_dtype(input)} + + self.attrs = { + 'strides': self.strides, + 'paddings': self.paddings, + 'ksize': self.ksize, + 'pooling_type': self.pool_type, + 'global_pooling': self.global_pool, + 'use_cudnn': self.use_cudnn, + 'use_mkldnn': self.use_mkldnn, + 'ceil_mode': self.ceil_mode, + 'data_format': self.data_format, + 'exclusive': self.exclusive, + 'adaptive': self.adaptive, + "padding_algorithm": self.padding_algorithm, + } + + self.outputs = {'Out': output} + + if self.use_cudnn: + self.python_api = pool2d_wrapper_use_cudnn + else: + self.python_api = pool2d_wrapper_not_use_cudnn + + def has_cudnn(self): + return core.is_compiled_with_cuda() and self.use_cudnn + + def test_check_output(self): + # TODO(wangzhongpu): support onednn op in dygraph mode + if self.has_cudnn(): + place = core.CUDAPlace(0) + self.check_output_with_place( + place, + atol=1e-5, + check_dygraph=(not self.use_mkldnn), + check_cinn=True, + check_pir=True, + check_pir_onednn=self.check_pir_onednn, + ) + else: + self.check_output( + check_dygraph=(not self.use_mkldnn), + check_pir=True, + check_pir_onednn=self.check_pir_onednn, + ) + + def test_check_grad(self): + if self.dtype == np.float16: + return + # TODO(wangzhongpu): support onednn op in dygraph mode + if self.has_cudnn() and self.pool_type != "max": + place = core.CUDAPlace(0) + self.check_grad_with_place( + place, + {'X'}, + 'Out', + check_dygraph=(not self.use_mkldnn), + check_cinn=True, + check_pir=True, + check_pir_onednn=self.check_pir_onednn, + ) + elif self.pool_type != "max": + self.check_grad( + {'X'}, + 'Out', + max_relative_error=0.07, + check_dygraph=(not self.use_mkldnn), + check_pir=True, + check_pir_onednn=self.check_pir_onednn, + ) + + def init_data_format(self): + self.data_format = "NCHW" + + def init_shape(self): + self.shape = [2, 3, 5, 5] + + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + + def init_paddings(self): + self.paddings = [0, 0] + self.padding_algorithm = "EXPLICIT" + + def init_kernel_type(self): + self.use_cudnn = False + + def init_data_type(self): + self.dtype = np.float32 if core.is_compiled_with_rocm() else np.float64 + + def init_pool_type(self): + self.pool_type = "avg" + self.pool2D_forward_naive = avg_pool2D_forward_naive + + def init_global_pool(self): + self.global_pool = True + + def init_ceil_mode(self): + self.ceil_mode = False + + def init_exclusive(self): + self.exclusive = True + + def init_adaptive(self): + self.adaptive = False + + +class TestPool2D_Op(TestPool2D_Op_Mixin, OpTest): + pass + + +class TestCase1(TestPool2D_Op): + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + + def init_paddings(self): + self.paddings = [0, 0] + + def init_pool_type(self): + self.pool_type = "avg" + self.pool2D_forward_naive = avg_pool2D_forward_naive + + def init_global_pool(self): + self.global_pool = False + + def init_shape(self): + self.shape = [2, 3, 7, 7] + + +class TestCase2(TestPool2D_Op): + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + + def init_paddings(self): + self.paddings = [1, 1] + + def init_pool_type(self): + self.pool_type = "avg" + self.pool2D_forward_naive = avg_pool2D_forward_naive + + def init_global_pool(self): + self.global_pool = False + + def init_shape(self): + self.shape = [2, 3, 7, 7] + + +class TestCase3(TestPool2D_Op): + def init_pool_type(self): + self.pool_type = "max" + self.pool2D_forward_naive = max_pool2D_forward_naive + + +class TestCase4(TestCase1): + def init_pool_type(self): + self.pool_type = "max" + self.pool2D_forward_naive = max_pool2D_forward_naive + + +class TestCase5(TestCase2): + def init_pool_type(self): + self.pool_type = "max" + self.pool2D_forward_naive = max_pool2D_forward_naive + + +# --------------------test pool2d cudnn-------------------- + + +def create_test_cudnn_class(parent): + @unittest.skipIf( + not core.is_compiled_with_cuda(), "core is not compiled with CUDA" + ) + class TestCUDNNCase(parent): + def init_kernel_type(self): + self.use_cudnn = True + + cls_name = "{}_{}".format(parent.__name__, "CUDNNOp") + TestCUDNNCase.__name__ = cls_name + globals()[cls_name] = TestCUDNNCase + + +create_test_cudnn_class(TestPool2D_Op) +create_test_cudnn_class(TestCase1) +create_test_cudnn_class(TestCase2) +create_test_cudnn_class(TestCase3) +create_test_cudnn_class(TestCase4) +create_test_cudnn_class(TestCase5) + +# --------------------test pool2d cudnn_fp16-------------------- + + +def create_test_cudnn_fp16_class(parent, check_grad=True): + @unittest.skipIf( + not core.is_compiled_with_cuda(), "core is not compiled with CUDA" + ) + class TestCUDNNFp16Case(parent): + def init_kernel_type(self): + self.use_cudnn = True + self.dtype = np.float16 + + def test_check_output(self): + # TODO(wangzhongpu): support onednn op in dygraph mode + if core.is_compiled_with_cuda(): + place = core.CUDAPlace(0) + if core.is_float16_supported(place): + self.check_output_with_place( + place, + check_dygraph=(not self.use_mkldnn), + check_cinn=True, + check_pir_onednn=self.check_pir_onednn, + ) + + def test_check_grad(self): + # TODO(wangzhongpu): support onednn op in dygraph mode + place = core.CUDAPlace(0) + if ( + core.is_float16_supported(place) + and self.pool_type != "max" + and check_grad + ): + self.check_grad_with_place( + place, + {'X'}, + 'Out', + check_dygraph=(not self.use_mkldnn), + check_cinn=True, + check_pir_onednn=self.check_pir_onednn, + ) + + cls_name = "{}_{}".format(parent.__name__, "CUDNNFp16Op") + TestCUDNNFp16Case.__name__ = cls_name + globals()[cls_name] = TestCUDNNFp16Case + + +def create_test_fp16_class(parent, check_grad=True): + @unittest.skipIf( + not core.is_compiled_with_cuda(), "core is not compiled with CUDA" + ) + class TestFp16Case(parent): + def init_kernel_type(self): + self.use_cudnn = False + self.dtype = np.float16 + + def test_check_output(self): + # TODO(wangzhongpu): support onednn op in dygraph mode + if core.is_compiled_with_cuda(): + place = core.CUDAPlace(0) + if core.is_float16_supported(place): + self.check_output_with_place( + place, + check_dygraph=(not self.use_mkldnn), + check_cinn=True, + check_pir_onednn=self.check_pir_onednn, + ) + + def test_check_grad(self): + # TODO(wangzhongpu): support onednn op in dygraph mode + place = core.CUDAPlace(0) + if ( + core.is_float16_supported(place) + and self.pool_type != "max" + and check_grad + ): + self.check_grad_with_place( + place, + {'X'}, + 'Out', + check_dygraph=(not self.use_mkldnn), + check_cinn=True, + check_pir_onednn=self.check_pir_onednn, + ) + + cls_name = "{}_{}".format(parent.__name__, "Fp16Op") + TestFp16Case.__name__ = cls_name + globals()[cls_name] = TestFp16Case + + +def create_test_bf16_class(parent, check_grad=True): + @unittest.skipIf( + not core.is_compiled_with_cuda(), "core is not compiled with CUDA" + ) + class TestBf16Case(parent): + def init_kernel_type(self): + self.use_cuda = True + self.dtype = np.uint16 + + def test_check_output(self): + if core.is_compiled_with_cuda(): + place = core.CUDAPlace(0) + self.check_output_with_place( + place, + check_dygraph=(not self.use_mkldnn), + check_cinn=True, + check_pir_onednn=self.check_pir_onednn, + ) + + def test_check_grad(self): + place = core.CUDAPlace(0) + if self.pool_type != "max" and check_grad: + self.check_grad_with_place( + place, + {'X'}, + 'Out', + check_dygraph=(not self.use_mkldnn), + check_cinn=True, + check_pir_onednn=self.check_pir_onednn, + ) + + cls_name = "{}_{}".format(parent.__name__, "Bf16Op") + TestBf16Case.__name__ = cls_name + globals()[cls_name] = TestBf16Case + + +create_test_cudnn_fp16_class(TestPool2D_Op) +create_test_cudnn_fp16_class(TestCase1) +create_test_cudnn_fp16_class(TestCase2) +create_test_cudnn_fp16_class(TestCase3) +create_test_cudnn_fp16_class(TestCase4) +create_test_cudnn_fp16_class(TestCase5) + +create_test_fp16_class(TestPool2D_Op) +create_test_fp16_class(TestCase1) +create_test_fp16_class(TestCase2) +create_test_fp16_class(TestCase3) +create_test_fp16_class(TestCase4) +create_test_fp16_class(TestCase5) + +create_test_bf16_class(TestPool2D_Op) +create_test_bf16_class(TestCase1) +create_test_bf16_class(TestCase2) +create_test_bf16_class(TestCase3) +create_test_bf16_class(TestCase4) +create_test_bf16_class(TestCase5) +# --------------------test pool2d use ceil mode-------------------- + + +def create_test_cudnn_use_ceil_class(parent): + @unittest.skipIf( + not core.is_compiled_with_cuda(), "core is not compiled with CUDA" + ) + class TestPool2DUseCeilCase(parent): + def init_kernel_type(self): + self.use_cudnn = True + + def init_ceil_mode(self): + self.ceil_mode = True + + cls_name = "{}_{}".format(parent.__name__, "CUDNNOpCeilMode") + TestPool2DUseCeilCase.__name__ = cls_name + globals()[cls_name] = TestPool2DUseCeilCase + + +create_test_cudnn_use_ceil_class(TestPool2D_Op) +create_test_cudnn_use_ceil_class(TestCase1) + + +def create_test_use_ceil_class(parent): + class TestPool2DUseCeilCase(parent): + def init_ceil_mode(self): + self.ceil_mode = True + + cls_name = "{}_{}".format(parent.__name__, "CeilModeCast") + TestPool2DUseCeilCase.__name__ = cls_name + globals()[cls_name] = TestPool2DUseCeilCase + + +create_test_use_ceil_class(TestCase1) +create_test_use_ceil_class(TestCase2) + + +class TestAvgInclude(TestCase2): + def init_exclusive(self): + self.exclusive = False + + +class TestCUDNNAvgInclude(TestCase2): + def init_kernel_type(self): + self.use_cudnn = True + + def init_exclusive(self): + self.exclusive = False + + +class TestAvgPoolAdaptive(TestCase1): + def init_adaptive(self): + self.adaptive = True + + +class TestAvgPoolAdaptiveAsyOutSize(TestCase1): + def init_adaptive(self): + self.adaptive = True + + def init_shape(self): + self.shape = [8, 3, 6, 6] + + def init_test_case(self): + self.ksize = [2, 3] + self.strides = [1, 1] + self.paddings = [0, 0, 0, 0] + + +# -------test pool2d with asymmetric padding----- + + +class TestPool2D_AsyPadding(TestPool2D_Op): + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + self.paddings = [1, 0, 1, 2] + + def init_shape(self): + self.shape = [2, 3, 5, 5] + + +class TestCase1_AsyPadding(TestCase1): + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + self.paddings = [1, 0, 1, 0] + + def init_shape(self): + self.shape = [2, 3, 7, 7] + + +class TestCase2_AsyPadding(TestCase2): + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + self.paddings = [1, 2, 1, 2] + + def init_shape(self): + self.shape = [2, 3, 7, 7] + + +class TestCase3_AsyPadding(TestCase3): + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + self.paddings = [1, 0, 1, 2] + + def init_shape(self): + self.shape = [2, 3, 5, 5] + + +class TestCase4_AsyPadding(TestCase4): + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + self.paddings = [1, 0, 1, 0] + + def init_shape(self): + self.shape = [2, 3, 7, 7] + + +class TestCase5_AsyPadding(TestCase5): + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + self.paddings = [2, 2, 1, 2] + + def init_shape(self): + self.shape = [2, 3, 7, 7] + + +create_test_cudnn_class(TestPool2D_AsyPadding) +create_test_cudnn_class(TestCase1_AsyPadding) +create_test_cudnn_class(TestCase2_AsyPadding) +create_test_cudnn_class(TestCase3_AsyPadding) +create_test_cudnn_class(TestCase4_AsyPadding) +create_test_cudnn_class(TestCase5_AsyPadding) + +create_test_cudnn_fp16_class(TestPool2D_AsyPadding) +create_test_cudnn_fp16_class(TestCase1_AsyPadding) +create_test_cudnn_fp16_class(TestCase2_AsyPadding) +create_test_cudnn_fp16_class(TestCase3_AsyPadding) +create_test_cudnn_fp16_class(TestCase4_AsyPadding) +create_test_cudnn_fp16_class(TestCase5_AsyPadding) + +create_test_fp16_class(TestPool2D_AsyPadding) +create_test_fp16_class(TestCase1_AsyPadding) +create_test_fp16_class(TestCase2_AsyPadding) +create_test_fp16_class(TestCase3_AsyPadding) +create_test_fp16_class(TestCase4_AsyPadding) +create_test_fp16_class(TestCase5_AsyPadding) + +create_test_bf16_class(TestPool2D_AsyPadding) +create_test_bf16_class(TestCase1_AsyPadding) +create_test_bf16_class(TestCase2_AsyPadding) +create_test_bf16_class(TestCase3_AsyPadding) +create_test_bf16_class(TestCase4_AsyPadding) +create_test_bf16_class(TestCase5_AsyPadding) + +create_test_cudnn_use_ceil_class(TestPool2D_AsyPadding) +create_test_cudnn_use_ceil_class(TestCase1_AsyPadding) + +create_test_use_ceil_class(TestCase1_AsyPadding) +create_test_use_ceil_class(TestCase2_AsyPadding) + + +class TestAvgInclude_AsyPadding(TestCase2): + def init_exclusive(self): + self.exclusive = False + + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + self.paddings = [1, 2, 1, 2] + + def init_shape(self): + self.shape = [2, 3, 7, 7] + + +class TestCUDNNAvgInclude_AsyPadding(TestCase2): + def init_kernel_type(self): + self.use_cudnn = True + + def init_exclusive(self): + self.exclusive = False + + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + self.paddings = [2, 1, 1, 1] + + def init_shape(self): + self.shape = [2, 3, 7, 7] + + +class TestAvgPoolAdaptive_AsyPadding(TestCase1): + def init_adaptive(self): + self.adaptive = True + + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 1] + self.paddings = [1, 1, 0, 2] + + def init_shape(self): + self.shape = [2, 3, 7, 7] + + +# ----------- test channel_last -------------- +class TestPool2D_channel_last(TestPool2D_Op): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 5, 5, 3] + + +class TestCase1_channel_last(TestCase1): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +class TestCase2_channel_last(TestCase2): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +class TestCase3_channel_last(TestCase3): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 5, 5, 3] + + +class TestCase4_channel_last(TestCase4): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +class TestCase5_channel_last(TestCase5): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +create_test_cudnn_class(TestPool2D_channel_last) +create_test_cudnn_class(TestCase1_channel_last) +create_test_cudnn_class(TestCase2_channel_last) +create_test_cudnn_class(TestCase3_channel_last) +create_test_cudnn_class(TestCase4_channel_last) +create_test_cudnn_class(TestCase5_channel_last) + +create_test_cudnn_fp16_class(TestPool2D_channel_last) +create_test_cudnn_fp16_class(TestCase1_channel_last) +create_test_cudnn_fp16_class(TestCase2_channel_last) +create_test_cudnn_fp16_class(TestCase3_channel_last) +create_test_cudnn_fp16_class(TestCase4_channel_last) +create_test_cudnn_fp16_class(TestCase5_channel_last) + +create_test_fp16_class(TestPool2D_channel_last) +create_test_fp16_class(TestCase1_channel_last) +create_test_fp16_class(TestCase2_channel_last) +create_test_fp16_class(TestCase3_channel_last) +create_test_fp16_class(TestCase4_channel_last) +create_test_fp16_class(TestCase5_channel_last) + +create_test_bf16_class(TestPool2D_channel_last) +create_test_bf16_class(TestCase1_channel_last) +create_test_bf16_class(TestCase2_channel_last) +create_test_bf16_class(TestCase3_channel_last) +create_test_bf16_class(TestCase4_channel_last) +create_test_bf16_class(TestCase5_channel_last) + +create_test_cudnn_use_ceil_class(TestPool2D_channel_last) +create_test_cudnn_use_ceil_class(TestCase1_channel_last) + +create_test_use_ceil_class(TestCase1_channel_last) +create_test_use_ceil_class(TestCase2_channel_last) + + +class TestCase5_Max(TestCase2): + def init_pool_type(self): + self.pool_type = "max" + + def test_check_grad(self): + if self.dtype == np.float16: + return + if self.has_cudnn() and self.pool_type == "max": + place = core.CUDAPlace(0) + self.check_grad_with_place( + place, + {'X'}, + 'Out', + max_relative_error=1.00, + check_cinn=True, + check_pir=True, + check_pir_onednn=self.check_pir_onednn, + ) + elif self.pool_type == "max": + self.check_grad( + {'X'}, + 'Out', + max_relative_error=1.00, + check_cinn=True, + check_pir=True, + check_pir_onednn=self.check_pir_onednn, + ) + + +class TestCase5_channel_last_Max(TestCase5_Max): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +create_test_cudnn_class(TestCase5_Max) +create_test_cudnn_class(TestCase5_channel_last_Max) + + +class TestAvgInclude_channel_last(TestCase2_channel_last): + def init_exclusive(self): + self.exclusive = False + + +class TestCUDNNAvgInclude_channel_last(TestCase2_channel_last): + def init_kernel_type(self): + self.use_cudnn = True + + def init_exclusive(self): + self.exclusive = False + + +class TestAvgPoolAdaptive_channel_last(TestCase1_channel_last): + def init_adaptive(self): + self.adaptive = True + + +class TestPool2D_AsyPadding_channel_last(TestPool2D_AsyPadding): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 5, 5, 3] + + +class TestCase1_AsyPadding_channel_last(TestCase1_AsyPadding): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +class TestCase2_AsyPadding_channel_last(TestCase2_AsyPadding): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +class TestCase3_AsyPadding_channel_last(TestCase3_AsyPadding): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 5, 5, 3] + + +class TestCase4_AsyPadding_channel_last(TestCase4_AsyPadding): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +class TestCase5_AsyPadding_channel_last(TestCase5_AsyPadding): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +create_test_cudnn_class(TestPool2D_AsyPadding_channel_last) +create_test_cudnn_class(TestCase1_AsyPadding_channel_last) +create_test_cudnn_class(TestCase2_AsyPadding_channel_last) +create_test_cudnn_class(TestCase3_AsyPadding_channel_last) +create_test_cudnn_class(TestCase4_AsyPadding_channel_last) +create_test_cudnn_class(TestCase5_AsyPadding_channel_last) + +create_test_cudnn_fp16_class(TestPool2D_AsyPadding_channel_last) +create_test_cudnn_fp16_class(TestCase1_AsyPadding_channel_last) +create_test_cudnn_fp16_class(TestCase2_AsyPadding_channel_last) +create_test_cudnn_fp16_class(TestCase3_AsyPadding_channel_last) +create_test_cudnn_fp16_class(TestCase4_AsyPadding_channel_last) +create_test_cudnn_fp16_class(TestCase5_AsyPadding_channel_last) + +create_test_fp16_class(TestPool2D_AsyPadding_channel_last) +create_test_fp16_class(TestCase1_AsyPadding_channel_last) +create_test_fp16_class(TestCase2_AsyPadding_channel_last) +create_test_fp16_class(TestCase3_AsyPadding_channel_last) +create_test_fp16_class(TestCase4_AsyPadding_channel_last) +create_test_fp16_class(TestCase5_AsyPadding_channel_last) + +create_test_bf16_class(TestPool2D_AsyPadding_channel_last) +create_test_bf16_class(TestCase1_AsyPadding_channel_last) +create_test_bf16_class(TestCase2_AsyPadding_channel_last) +create_test_bf16_class(TestCase3_AsyPadding_channel_last) +create_test_bf16_class(TestCase4_AsyPadding_channel_last) +create_test_bf16_class(TestCase5_AsyPadding_channel_last) + +create_test_cudnn_use_ceil_class(TestPool2D_AsyPadding_channel_last) +create_test_cudnn_use_ceil_class(TestCase1_AsyPadding_channel_last) + +create_test_use_ceil_class(TestCase1_AsyPadding_channel_last) +create_test_use_ceil_class(TestCase2_AsyPadding_channel_last) + + +class TestAvgInclude_AsyPadding_channel_last(TestAvgInclude_AsyPadding): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +class TestCUDNNAvgInclude_AsyPadding_channel_last( + TestCUDNNAvgInclude_AsyPadding +): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +class TestAvgPoolAdaptive_AsyPadding_channel_last( + TestAvgPoolAdaptive_AsyPadding +): + def init_data_format(self): + self.data_format = "NHWC" + + def init_shape(self): + self.shape = [2, 7, 7, 3] + + +# test paddings: SAME VALID + + +def create_test_padding_SAME_class(parent): + class TestPaddingSMAECase(parent): + def init_paddings(self): + self.paddings = [0, 0] + self.padding_algorithm = "SAME" + + cls_name = "{}_{}".format(parent.__name__, "PaddingSAMEOp") + TestPaddingSMAECase.__name__ = cls_name + globals()[cls_name] = TestPaddingSMAECase + + +create_test_padding_SAME_class(TestPool2D_Op) +create_test_padding_SAME_class(TestCase1) +create_test_padding_SAME_class(TestCase2) +create_test_padding_SAME_class(TestCase3) +create_test_padding_SAME_class(TestCase4) +create_test_padding_SAME_class(TestCase5) + +create_test_padding_SAME_class(TestPool2D_channel_last) +create_test_padding_SAME_class(TestCase1_channel_last) +create_test_padding_SAME_class(TestCase2_channel_last) +create_test_padding_SAME_class(TestCase3_channel_last) +create_test_padding_SAME_class(TestCase4_channel_last) +create_test_padding_SAME_class(TestCase5_channel_last) + + +def create_test_cudnn_padding_SAME_class(parent): + @unittest.skipIf( + not core.is_compiled_with_cuda(), "core is not compiled with CUDA" + ) + class TestCUDNNPaddingSMAECase(parent): + def init_kernel_type(self): + self.use_cudnn = True + + def init_paddings(self): + self.paddings = [1, 1] + self.padding_algorithm = "SAME" + + cls_name = "{}_{}".format(parent.__name__, "CudnnPaddingSAMEOp") + TestCUDNNPaddingSMAECase.__name__ = cls_name + globals()[cls_name] = TestCUDNNPaddingSMAECase + + +create_test_cudnn_padding_SAME_class(TestPool2D_Op) +create_test_cudnn_padding_SAME_class(TestCase1) +create_test_cudnn_padding_SAME_class(TestCase2) +create_test_cudnn_padding_SAME_class(TestCase3) +create_test_cudnn_padding_SAME_class(TestCase4) +create_test_cudnn_padding_SAME_class(TestCase5) + +create_test_cudnn_padding_SAME_class(TestPool2D_channel_last) +create_test_cudnn_padding_SAME_class(TestCase1_channel_last) +create_test_cudnn_padding_SAME_class(TestCase2_channel_last) +create_test_cudnn_padding_SAME_class(TestCase3_channel_last) +create_test_cudnn_padding_SAME_class(TestCase4_channel_last) +create_test_cudnn_padding_SAME_class(TestCase5_channel_last) + + +def create_test_padding_VALID_class(parent): + class TestPaddingVALIDCase(parent): + def init_paddings(self): + self.paddings = [1, 1] + self.padding_algorithm = "VALID" + + cls_name = "{}_{}".format(parent.__name__, "PaddingVALIDOp") + TestPaddingVALIDCase.__name__ = cls_name + globals()[cls_name] = TestPaddingVALIDCase + + +create_test_padding_VALID_class(TestPool2D_Op) +create_test_padding_VALID_class(TestCase1) +create_test_padding_VALID_class(TestCase2) +create_test_padding_VALID_class(TestCase3) +create_test_padding_VALID_class(TestCase4) +create_test_padding_VALID_class(TestCase5) + +create_test_padding_VALID_class(TestPool2D_channel_last) +create_test_padding_VALID_class(TestCase1_channel_last) +create_test_padding_VALID_class(TestCase2_channel_last) +create_test_padding_VALID_class(TestCase3_channel_last) +create_test_padding_VALID_class(TestCase4_channel_last) +create_test_padding_VALID_class(TestCase5_channel_last) + + +def create_test_cudnn_padding_VALID_class(parent): + @unittest.skipIf( + not core.is_compiled_with_cuda(), "core is not compiled with CUDA" + ) + class TestCUDNNPaddingVALIDCase(parent): + def init_kernel_type(self): + self.use_cudnn = True + + def init_paddings(self): + self.paddings = [1, 1] + self.padding_algorithm = "VALID" + + cls_name = "{}_{}".format(parent.__name__, "CudnnPaddingVALIDOp") + TestCUDNNPaddingVALIDCase.__name__ = cls_name + globals()[cls_name] = TestCUDNNPaddingVALIDCase + + +create_test_cudnn_padding_VALID_class(TestPool2D_Op) +create_test_cudnn_padding_VALID_class(TestCase1) +create_test_cudnn_padding_VALID_class(TestCase2) +create_test_cudnn_padding_VALID_class(TestCase3) +create_test_cudnn_padding_VALID_class(TestCase4) +create_test_cudnn_padding_VALID_class(TestCase5) + +create_test_cudnn_padding_VALID_class(TestPool2D_channel_last) +create_test_cudnn_padding_VALID_class(TestCase1_channel_last) +create_test_cudnn_padding_VALID_class(TestCase2_channel_last) +create_test_cudnn_padding_VALID_class(TestCase3_channel_last) +create_test_cudnn_padding_VALID_class(TestCase4_channel_last) +create_test_cudnn_padding_VALID_class(TestCase5_channel_last) + + +class TestCase1_strides(TestCase1): + def init_test_case(self): + self.ksize = [3, 3] + self.strides = [1, 2] + + def init_shape(self): + self.shape = [2, 3, 4, 5] + + +create_test_cudnn_class(TestCase1_strides) +create_test_padding_SAME_class(TestCase1_strides) +create_test_cudnn_padding_SAME_class(TestCase1_strides) + + +if __name__ == '__main__': + unittest.main() From 081e6228d0f5a79c6784b88ada39b1c8eb3a3279 Mon Sep 17 00:00:00 2001 From: Wang Huan Date: Tue, 14 May 2024 06:09:45 +0000 Subject: [PATCH 3/6] refine --- test/deprecated/legacy_test/CMakeLists.txt | 5 ----- test/legacy_test/CMakeLists.txt | 7 ++++++- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/test/deprecated/legacy_test/CMakeLists.txt b/test/deprecated/legacy_test/CMakeLists.txt index 4968e979a137ad..6f07ebe52c0c35 100644 --- a/test/deprecated/legacy_test/CMakeLists.txt +++ b/test/deprecated/legacy_test/CMakeLists.txt @@ -723,7 +723,6 @@ set_tests_properties(test_reduce_op PROPERTIES TIMEOUT 500) set_tests_properties(test_conv_nn_grad PROPERTIES TIMEOUT 220) set_tests_properties(test_program_prune_backward PROPERTIES TIMEOUT 120) set_tests_properties(test_imperative_optimizer_v2 PROPERTIES TIMEOUT 250) -set_tests_properties(test_pool2d_op PROPERTIES TIMEOUT 120) set_tests_properties(test_transpose_op PROPERTIES TIMEOUT 120) set_tests_properties(test_bilinear_interp_op PROPERTIES TIMEOUT 120) set_tests_properties(test_decoupled_py_reader PROPERTIES TIMEOUT 120) @@ -760,8 +759,6 @@ if(APPLE) PROPERTIES TIMEOUT 300) endif() -set_tests_properties(test_inplace_addto_strategy PROPERTIES TIMEOUT 120) - set(TEST_CINN_OPS test_softmax_op test_expand_v2_op @@ -791,7 +788,6 @@ set(TEST_CINN_OPS test_flip test_triangular_solve_op test_scatter_nd_op - test_pool2d_op test_instance_norm_op test_cumsum_op test_split_op @@ -891,6 +887,5 @@ set_tests_properties(test_linalg_matrix_exp PROPERTIES TIMEOUT 120) set_pir_tests_properties() set_tests_properties(test_fractional_max_pool2d_op PROPERTIES TIMEOUT 120) -set_tests_properties(test_fractional_max_pool3d_op PROPERTIES TIMEOUT 120) set_tests_properties(test_reduce_as_op PROPERTIES TIMEOUT 30) diff --git a/test/legacy_test/CMakeLists.txt b/test/legacy_test/CMakeLists.txt index 767775eb17db41..ba70f3d2efad2f 100644 --- a/test/legacy_test/CMakeLists.txt +++ b/test/legacy_test/CMakeLists.txt @@ -857,6 +857,10 @@ set_tests_properties(test_dataset_uci_housing PROPERTIES TIMEOUT 120) set_tests_properties(test_dataset_imdb PROPERTIES TIMEOUT 300) set_tests_properties(test_callback_wandb PROPERTIES TIMEOUT 60) set_tests_properties(test_jit_save_load PROPERTIES TIMEOUT 100) +set_tests_properties(test_pool2d_op PROPERTIES TIMEOUT 120) +set_tests_properties(test_fractional_max_pool3d_op PROPERTIES TIMEOUT 120) +set_tests_properties(test_inplace_addto_strategy PROPERTIES TIMEOUT 120) + if(WITH_COVERAGE) set_tests_properties(test_hapi_hub PROPERTIES TIMEOUT 300) endif() @@ -946,7 +950,8 @@ set(TEST_CINN_OPS test_elementwise_min_op test_take_along_axis_op test_strided_slice_op - test_pad_op) + test_pad_op + test_pool2d_op) foreach(TEST_CINN_OP ${TEST_CINN_OPS}) if(WITH_CINN) From f23d5d36ded2bab4b68a3e510d1f12a4fa503bbf Mon Sep 17 00:00:00 2001 From: Wang Huan Date: Tue, 14 May 2024 08:19:23 +0000 Subject: [PATCH 4/6] refine --- test/deprecated/legacy_test/CMakeLists.txt | 3 +++ test/legacy_test/CMakeLists.txt | 1 - 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/test/deprecated/legacy_test/CMakeLists.txt b/test/deprecated/legacy_test/CMakeLists.txt index 6f07ebe52c0c35..fd8f03137855d2 100644 --- a/test/deprecated/legacy_test/CMakeLists.txt +++ b/test/deprecated/legacy_test/CMakeLists.txt @@ -759,6 +759,9 @@ if(APPLE) PROPERTIES TIMEOUT 300) endif() +set_tests_properties(test_inplace_addto_strategy_deprecated PROPERTIES TIMEOUT + 120) + set(TEST_CINN_OPS test_softmax_op test_expand_v2_op diff --git a/test/legacy_test/CMakeLists.txt b/test/legacy_test/CMakeLists.txt index ba70f3d2efad2f..957d995529e74c 100644 --- a/test/legacy_test/CMakeLists.txt +++ b/test/legacy_test/CMakeLists.txt @@ -859,7 +859,6 @@ set_tests_properties(test_callback_wandb PROPERTIES TIMEOUT 60) set_tests_properties(test_jit_save_load PROPERTIES TIMEOUT 100) set_tests_properties(test_pool2d_op PROPERTIES TIMEOUT 120) set_tests_properties(test_fractional_max_pool3d_op PROPERTIES TIMEOUT 120) -set_tests_properties(test_inplace_addto_strategy PROPERTIES TIMEOUT 120) if(WITH_COVERAGE) set_tests_properties(test_hapi_hub PROPERTIES TIMEOUT 300) From 56cda2e467a5040d66abb41cf002c9368eb36f26 Mon Sep 17 00:00:00 2001 From: Wang Huan Date: Wed, 15 May 2024 02:34:41 +0000 Subject: [PATCH 5/6] refine --- .../legacy_test/test_elementwise_gradient_op_deprecated.py | 3 +++ .../legacy_test/test_inplace_addto_strategy_deprecated.py | 3 ++- test/deprecated/legacy_test/test_set_bool_attr_deprecated.py | 1 + 3 files changed, 6 insertions(+), 1 deletion(-) diff --git a/test/deprecated/legacy_test/test_elementwise_gradient_op_deprecated.py b/test/deprecated/legacy_test/test_elementwise_gradient_op_deprecated.py index e06ac89200a7cd..e07ddec9198550 100644 --- a/test/deprecated/legacy_test/test_elementwise_gradient_op_deprecated.py +++ b/test/deprecated/legacy_test/test_elementwise_gradient_op_deprecated.py @@ -16,6 +16,7 @@ import numpy as np +import paddle from paddle import base from paddle.base import core @@ -104,6 +105,7 @@ def test_with_place(place): test_with_place(place) def test_check_forward_backward_with_scale_and_bias(self): + paddle.enable_static() np.random.seed(123) self.x = np.random.random((4, 32, 220, 220)).astype(np.float32) self.y = np.random.random(32).astype(np.float32) @@ -113,4 +115,5 @@ def test_check_forward_backward_with_scale_and_bias(self): if __name__ == '__main__': + paddle.enable_static() unittest.main() diff --git a/test/deprecated/legacy_test/test_inplace_addto_strategy_deprecated.py b/test/deprecated/legacy_test/test_inplace_addto_strategy_deprecated.py index b0f21fbbba709e..e34bd71fa59c17 100644 --- a/test/deprecated/legacy_test/test_inplace_addto_strategy_deprecated.py +++ b/test/deprecated/legacy_test/test_inplace_addto_strategy_deprecated.py @@ -113,12 +113,13 @@ def run_program(enable_addto): np.testing.assert_array_equal(res1, res2) def test_nchw(self): + paddle.enable_static() self.check_result() def test_nhwc(self): + paddle.enable_static() self.check_result("NHWC") if __name__ == "__main__": - paddle.enable_static() unittest.main() diff --git a/test/deprecated/legacy_test/test_set_bool_attr_deprecated.py b/test/deprecated/legacy_test/test_set_bool_attr_deprecated.py index 1752a68f09e089..3e2d91a8262027 100644 --- a/test/deprecated/legacy_test/test_set_bool_attr_deprecated.py +++ b/test/deprecated/legacy_test/test_set_bool_attr_deprecated.py @@ -21,6 +21,7 @@ class TestAttrSet(unittest.TestCase): def test_set_bool_attr(self): + paddle.enable_static() x = paddle.static.data( name='x', shape=[-1, 3, 7, 3, 7], dtype='float32' ) From 544f920c801478ffa46649829a7ab9ef157f12ad Mon Sep 17 00:00:00 2001 From: Wang Huan Date: Thu, 16 May 2024 01:45:52 +0000 Subject: [PATCH 6/6] refine --- test/deprecated/legacy_test/CMakeLists.txt | 2 -- 1 file changed, 2 deletions(-) diff --git a/test/deprecated/legacy_test/CMakeLists.txt b/test/deprecated/legacy_test/CMakeLists.txt index 0c8d84368bae58..c8559b7aaf78d9 100644 --- a/test/deprecated/legacy_test/CMakeLists.txt +++ b/test/deprecated/legacy_test/CMakeLists.txt @@ -722,8 +722,6 @@ set_tests_properties(test_reduce_op PROPERTIES TIMEOUT 500) set_tests_properties(test_conv_nn_grad PROPERTIES TIMEOUT 220) set_tests_properties(test_program_prune_backward PROPERTIES TIMEOUT 120) set_tests_properties(test_imperative_optimizer_v2 PROPERTIES TIMEOUT 250) -set_tests_properties(test_transpose_op PROPERTIES TIMEOUT 120) -set_tests_properties(test_pool2d_op PROPERTIES TIMEOUT 120) set_tests_properties(test_bilinear_interp_op PROPERTIES TIMEOUT 120) set_tests_properties(test_decoupled_py_reader PROPERTIES TIMEOUT 120) set_tests_properties(test_fuse_bn_act_pass PROPERTIES TIMEOUT 120)