diff --git a/tests/ttnn/unit_tests/operations/test_upsample.py b/tests/ttnn/unit_tests/operations/test_upsample.py index d2795db29e8..7405bfda43e 100644 --- a/tests/ttnn/unit_tests/operations/test_upsample.py +++ b/tests/ttnn/unit_tests/operations/test_upsample.py @@ -119,6 +119,7 @@ def test_upsample_single_core(device, input_shapes, scale_h, scale_w): [1, 64, 132, 19], ], ) +@pytest.mark.parametrize("device_params", [{"l1_small_size": 24576}], indirect=True) @pytest.mark.parametrize("scale_h", [2, 3]) @pytest.mark.parametrize("scale_w", [2, 3]) @pytest.mark.parametrize("shard_strategy", [ttnn.ShardStrategy.HEIGHT, ttnn.ShardStrategy.BLOCK]) diff --git a/ttnn/cpp/ttnn/operations/pool/upsample/device/upsample_program_factory_multicore.cpp b/ttnn/cpp/ttnn/operations/pool/upsample/device/upsample_program_factory_multicore.cpp index 3165c85e97c..d8c11117032 100644 --- a/ttnn/cpp/ttnn/operations/pool/upsample/device/upsample_program_factory_multicore.cpp +++ b/ttnn/cpp/ttnn/operations/pool/upsample/device/upsample_program_factory_multicore.cpp @@ -4,15 +4,19 @@ #include -#include "buffers/buffer.hpp" -#include "buffers/buffer_constants.hpp" -#include "common/assert.hpp" -#include "common/core_coord.hpp" -#include "ttnn/tensor/host_buffer/functions.hpp" +#include + +#include "upsample_op.hpp" +#include "ttnn/operations/math.hpp" #include "tt_metal/host_api.hpp" +#include "tt_metal/common/constants.hpp" +#include "tt_metal/detail/util.hpp" #include "tt_metal/common/math.hpp" +#include "tt_metal/tt_stl/reflection.hpp" +#include "ttnn/tensor/host_buffer/functions.hpp" + using namespace tt::constants; using namespace tt::tt_metal; @@ -229,7 +233,7 @@ operation::ProgramWithCallbacks upsample_multi_core( ? ShardOrientation::COL_MAJOR : shard_spec.orientation; ShardSpec config_shard_spec(input.shard_spec().value().grid, shard_shape, config_tensor_shard_orientation, false); - MemoryConfig memory_config{input.memory_config().memory_layout, BufferType::L1, config_shard_spec}; + MemoryConfig memory_config{input.memory_config().memory_layout, BufferType::L1_SMALL, config_shard_spec}; auto config_tensor_device = config_tensor.to(device, memory_config); tt::tt_metal::detail::AddConfigBuffer(program, config_tensor_device.device_buffer());