From 91cb7336400acc857e20086a23692f99fe11be9c Mon Sep 17 00:00:00 2001 From: Gunes Bayir Date: Tue, 25 Jul 2023 17:00:33 +0100 Subject: Port Resize operator to CKW Use Compute Kernel Writer (CKW) to generate code for Resize operator in the Dynamic Fusion interface. Supports Nearest Neighbor and Bilinear interpolation methods. Resolves: COMPMID-6265 Change-Id: Ib0a5158bd4208123c84f6a1dc54f29d82fd55dcd Signed-off-by: Gunes Bayir Signed-off-by: Jakub Sujak Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/10174 Tested-by: Arm Jenkins Comments-Addressed: Arm Jenkins Benchmark: Arm Jenkins --- compute_kernel_writer/prototype/include/ckw/KernelWriter.h | 2 +- compute_kernel_writer/prototype/include/ckw/types/Functions.h | 4 +++- compute_kernel_writer/prototype/src/KernelWriter.cpp | 2 +- compute_kernel_writer/prototype/src/Prototype.h | 6 ++++++ 4 files changed, 11 insertions(+), 3 deletions(-) (limited to 'compute_kernel_writer') diff --git a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h index 72f85c78aa..fdb5fedc59 100644 --- a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h +++ b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h @@ -282,7 +282,7 @@ public: * @param[out] dst The tile to be written to. * @param[in] dim The global ID dimension. */ - void op_get_global_id(TileOperand &dst, int32_t dim); + void op_get_global_id(const TileOperand &dst, int32_t dim); // ============================================================================================= // Code generation diff --git a/compute_kernel_writer/prototype/include/ckw/types/Functions.h b/compute_kernel_writer/prototype/include/ckw/types/Functions.h index 2dd5ed0b3d..bc1f85c188 100644 --- a/compute_kernel_writer/prototype/include/ckw/types/Functions.h +++ b/compute_kernel_writer/prototype/include/ckw/types/Functions.h @@ -39,9 +39,10 @@ enum class UnaryFunction : int32_t Fabs = 0x0004, Log = 0x0006, Round = 0x0007, + Floor = 0x0008, // Misc - SizeOf = 0x0008, + SizeOf = 0x0009, }; enum class BinaryFunction : int32_t @@ -53,6 +54,7 @@ enum class BinaryFunction : int32_t enum class TernaryFunction : int32_t { Select = 0x0000, + Clamp = 0x0001, }; } // namespace ckw diff --git a/compute_kernel_writer/prototype/src/KernelWriter.cpp b/compute_kernel_writer/prototype/src/KernelWriter.cpp index f29cf12802..5c9a16ee33 100644 --- a/compute_kernel_writer/prototype/src/KernelWriter.cpp +++ b/compute_kernel_writer/prototype/src/KernelWriter.cpp @@ -341,7 +341,7 @@ void KernelWriter::op_for_loop(const TileOperand &var_name, BinaryOp cond_op, co // Misc // ================================================================================================= -void KernelWriter::op_get_global_id(TileOperand &dst, int32_t dim) +void KernelWriter::op_get_global_id(const TileOperand &dst, int32_t dim) { _impl->op_get_global_id(prototype::Operand(dst.name()), dim); } diff --git a/compute_kernel_writer/prototype/src/Prototype.h b/compute_kernel_writer/prototype/src/Prototype.h index 2b519471ac..88d6e898e4 100644 --- a/compute_kernel_writer/prototype/src/Prototype.h +++ b/compute_kernel_writer/prototype/src/Prototype.h @@ -3694,6 +3694,9 @@ public: case UnaryFunction::Round: _data->code += "round("; break; + case UnaryFunction::Floor: + _data->code += "floor("; + break; default: CKW_ASSERT_MSG(false, "Unexpected UnaryFunction used."); } @@ -3772,6 +3775,9 @@ public: case TernaryFunction::Select: _data->code += "select("; break; + case TernaryFunction::Clamp: + _data->code += "clamp("; + break; default: CKW_ASSERT_MSG(false, "Unexpected TernaryFunction used."); } -- cgit v1.2.1