commit 91cb7336400acc857e20086a23692f99fe11be9c
author Gunes Bayir <gunes.bayir@arm.com> Tue Jul 25 17:00:33 2023 +0100
committer Jakub Sujak <jakub.sujak@arm.com> Wed Aug 30 13:48:07 2023 +0000
tree 188c96c61a197ab3a21106d195be4a6a4f63ca6c
parent d0d8f2e61039826685aa076347eacce526e8c74b

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>

compute_kernel_writer/prototype/include/ckw/KernelWriter.h

diff --git a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h
index 72f85c7..fdb5fed 100644
--- a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h
+++ b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h
@@ -282,7 +282,7 @@
      * @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

compute_kernel_writer/prototype/include/ckw/types/Functions.h

diff --git a/compute_kernel_writer/prototype/include/ckw/types/Functions.h b/compute_kernel_writer/prototype/include/ckw/types/Functions.h
index 2dd5ed0..bc1f85c 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 @@
     Fabs           = 0x0004,
     Log            = 0x0006,
     Round          = 0x0007,
+    Floor          = 0x0008,
 
     // Misc
-    SizeOf = 0x0008,
+    SizeOf = 0x0009,
 };
 
 enum class BinaryFunction : int32_t
@@ -53,6 +54,7 @@
 enum class TernaryFunction : int32_t
 {
     Select = 0x0000,
+    Clamp  = 0x0001,
 };
 
 } // namespace ckw

compute_kernel_writer/prototype/src/KernelWriter.cpp

diff --git a/compute_kernel_writer/prototype/src/KernelWriter.cpp b/compute_kernel_writer/prototype/src/KernelWriter.cpp
index f29cf12..5c9a16e 100644
--- a/compute_kernel_writer/prototype/src/KernelWriter.cpp
+++ b/compute_kernel_writer/prototype/src/KernelWriter.cpp
@@ -341,7 +341,7 @@
 // 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);
 }

compute_kernel_writer/prototype/src/Prototype.h

diff --git a/compute_kernel_writer/prototype/src/Prototype.h b/compute_kernel_writer/prototype/src/Prototype.h
index 2b51947..88d6e89 100644
--- a/compute_kernel_writer/prototype/src/Prototype.h
+++ b/compute_kernel_writer/prototype/src/Prototype.h
@@ -3694,6 +3694,9 @@
                 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 @@
                 case TernaryFunction::Select:
                     _data->code += "select(";
                     break;
+                case TernaryFunction::Clamp:
+                    _data->code += "clamp(";
+                    break;
                 default:
                     CKW_ASSERT_MSG(false, "Unexpected TernaryFunction used.");
             }