commit 16c5697085c256c19fb8ba4bef6188d61f30a88b
author Gunes Bayir <gunes.bayir@arm.com> Mon Mar 28 21:32:33 2022 +0100
committer SiCong Li <sicong.li@arm.com> Wed Apr 13 10:36:30 2022 +0000
tree 609bfe2082c939ff37bdf6ef37bc22fc071bd934
parent 5d606cccaabdfc435734c9fb51e11f14f3724a23

Add DirectConvolution2D kernel component for dynamic fusion

Resolves: COMPMID-5156

Change-Id: I438da924cb80d3bce72106b06ca7181e0606bd01
Signed-off-by: Gunes Bayir <gunes.bayir@arm.com>
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/7399
Reviewed-by: SiCong Li <sicong.li@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>

src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/Common.h

diff --git a/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/Common.h b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/Common.h
index 4c720ea..e24c742 100644
--- a/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/Common.h
+++ b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/Common.h
@@ -31,6 +31,7 @@
 #include "arm_compute/core/Error.h"
 #include "arm_compute/core/GPUTarget.h"
 #include "src/core/common/Macros.h"
+#include "support/Requires.h"
 #include "support/StringSupport.h"
 
 #include "src/core/experimental/dynamic_fusion/ClKernelBuildingAPI.h"
@@ -198,8 +199,9 @@
         {
         }
 
-        TagVal(ComponentID id)
-            : value{ std::to_string(id) }
+        template <typename T, ARM_COMPUTE_REQUIRES_TA(std::is_integral<T>::value)>
+        TagVal(T val)
+            : value{ support::cpp11::to_string(val) }
         {
         }
 
@@ -208,6 +210,16 @@
         {
         }
 
+        TagVal(const char *val)
+            : value{ std::string(val) }
+        {
+        }
+
+        TagVal(const DataType &data_type)
+            : value{ get_cl_type_from_data_type(data_type) }
+        {
+        }
+
         std::string value{};
     };
     using TagLUT = std::unordered_map<Tag, TagVal>; // Used to instantiating a code template / replacing tags
@@ -633,21 +645,36 @@
         std::string code;
         switch(var.desc.tensor_arg_type)
         {
+            case TensorArgType::Vector:
+            {
+                code += "\n    VECTOR_DECLARATION(" + var.uniq_name + ")";
+                break;
+            }
             case TensorArgType::Image:
             {
-                code += "IMAGE_DECLARATION(" + var.uniq_name + ")";
+                code += "\n    IMAGE_DECLARATION(" + var.uniq_name + ")";
                 break;
             }
             case TensorArgType::Image_3D:
             {
-                code += "IMAGE_DECLARATION(" + var.uniq_name + "),\n";
-                code += "uint " + var.uniq_name + "_stride_z";
+                code += "\n    IMAGE_DECLARATION(" + var.uniq_name + "),";
+                code += "\n    uint " + var.uniq_name + "_stride_z";
                 break;
             }
             case TensorArgType::Image_3D_Export_To_ClImage2D:
             {
-                code += "__read_only image2d_t " + var.uniq_name + "_img,\n";
-                code += "uint " + var.uniq_name + "_stride_z,\n";
+                code += "\n    __read_only image2d_t " + var.uniq_name + "_img,";
+                code += "\n    uint " + var.uniq_name + "_stride_z";
+                break;
+            }
+            case TensorArgType::Tensor_4D_t_Buffer:
+            {
+                code += "\n    TENSOR4D_T(" + var.uniq_name + ", BUFFER)";
+                break;
+            }
+            case TensorArgType::Tensor_4D_t_Image:
+            {
+                code += "\n    TENSOR4D_T(" + var.uniq_name + ", IMAGE)";
                 break;
             }
             default:
@@ -664,7 +691,7 @@
 
         for(const auto &arg : argument_list)
         {
-            code += "\n    " + generate_argument_declaration(arg) + ",";
+            code += generate_argument_declaration(arg) + ",";
         }
 
         code[code.length() - 1] = ')';
@@ -674,7 +701,8 @@
 
     std::string generate_global_section() const
     {
-        std::string code = "    uint g_x = get_global_id(0);\n";
+        std::string code = "";
+        code += "    uint g_x = get_global_id(0);\n";
         code += "    uint g_y = get_global_id(1);\n";
         code += "    uint g_z = get_global_id(2);\n\n";