diff options
| author | Gunes Bayir <gunes.bayir@arm.com> | 2023-07-25 17:00:33 +0100 |
|---|---|---|
| committer | Jakub Sujak <jakub.sujak@arm.com> | 2023-08-30 13:48:07 +0000 |
| commit | 91cb7336400acc857e20086a23692f99fe11be9c (patch) | |
| tree | 188c96c61a197ab3a21106d195be4a6a4f63ca6c /compute_kernel_writer | |
| parent | d0d8f2e61039826685aa076347eacce526e8c74b (diff) | |
| download | ComputeLibrary-91cb7336400acc857e20086a23692f99fe11be9c.tar.gz | |
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 <gunes.bayir@arm.com>
Signed-off-by: Jakub Sujak <jakub.sujak@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/10174
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Benchmark: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'compute_kernel_writer')
4 files changed, 11 insertions, 3 deletions
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."); } |