blob: e24c742fd7cae1d8656deb3bcc9c56fe720e580d [file] [log] [blame]
Giorgio Arena232c4522022-03-03 10:09:01 +00001/*
2 * Copyright (c) 2022 Arm Limited.
3 *
4 * SPDX-License-Identifier: MIT
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a copy
7 * of this software and associated documentation files (the "Software"), to
8 * deal in the Software without restriction, including without limitation the
9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10 * sell copies of the Software, and to permit persons to whom the Software is
11 * furnished to do so, subject to the following conditions:
12 *
13 * The above copyright notice and this permission notice shall be included in all
14 * copies or substantial portions of the Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22 * SOFTWARE.
23 */
24#if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
25
26#ifndef ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMMON_H
27#define ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMMON_H
28
29#include "arm_compute/core/CL/CLCompileContext.h"
Giorgio Arena892b70a2022-03-30 12:23:10 +010030#include "arm_compute/core/CL/CLKernelLibrary.h"
Giorgio Arena232c4522022-03-03 10:09:01 +000031#include "arm_compute/core/Error.h"
32#include "arm_compute/core/GPUTarget.h"
Gunes Bayir8a879832022-03-10 21:21:01 +000033#include "src/core/common/Macros.h"
Gunes Bayir16c56972022-03-28 21:32:33 +010034#include "support/Requires.h"
Giorgio Arenabd44caa2022-03-15 13:45:15 +000035#include "support/StringSupport.h"
Giorgio Arena232c4522022-03-03 10:09:01 +000036
37#include "src/core/experimental/dynamic_fusion/ClKernelBuildingAPI.h"
38
39#include <queue>
40#include <stack>
41#include <string>
42#include <unordered_set>
43
44namespace arm_compute
45{
46namespace experimental
47{
48namespace dynamic_fusion
49{
50/** We introduce the concept of *Shared Variables* in the context of kernel building.
51 * They are variables that can be accessed / shared among all the kernel components within a single kernel.
52 * For now we consider 2 groups of shared variables:
53 * Argument: The argument variables (parameters) of a kernel
54 * Automatic: The automatic variables declared inside a kernel
55 * All Shared Variables have the same kernel scope, and are thus visible to all kernel components
56*/
57
58enum class SharedVarIO
59{
60 Input,
61 Output
62};
63
64enum class SharedVarGroup
65{
66 Argument, // Parameters to a kernel function
67 Automatic // Automatic variables declared within the kernel body
68};
69
Gunes Bayir8a879832022-03-10 21:21:01 +000070/** Specifies a shared variable link for a component.
71 * It describes all the information that's available when a component is constructed / added:
Giorgio Arena232c4522022-03-03 10:09:01 +000072 * e.g. its linkage (via ArgumentID and io) and its group
73 * This is not shared variable on its own, but is used for instantiating a SharedVar when building the code
74 */
75struct SharedVarLink
76{
77 ArgumentID arg_id{ g_arg_placeholder };
78 SharedVarIO io{ SharedVarIO::Input };
79 SharedVarGroup group{ SharedVarGroup::Argument };
80 bool is_empty() const
81 {
82 return arg_id == g_arg_placeholder;
83 }
84};
85
86/** A table of all the variables used in the kernel / blueprint
87 * NOTE: the order they appear in the table is the order of their "declaration" in the component code, and is also their ID
88 * NOTE: the variables all have the scope of the full kernel function
89 */
90class SharedVarTable
91{
92public:
93 struct SharedVar
94 {
95 SharedVarGroup group;
96 std::string uniq_name; // Unique name, also the final variable name used in the built code
97 ClKernelArgRuntimeDescriptor desc; // Automatic variables can and should still be described using this struct
98 };
99
100 using Arguments = std::vector<SharedVar>;
101
102 /** @note: The order of insertion is important. There is one precondition:
103 * PRECOND: The components have been sorted topologically / is being traversed in topological order
104 * This ensures that all the consumer var links (Output, Automatic Links) can consume (return) the producer var links when they're referred
105 */
106 SharedVar add(SharedVarLink var_link, ClKernelArgRuntimeDescriptor runtime_desc, const std::string &name = "unnamed")
107 {
108 ARM_COMPUTE_ERROR_ON_MSG(var_link.is_empty(), "Non-empty SharedVarLink expected");
109 auto var_id = _num_var;
110 std::stringstream ss;
111 ss << name << "_" << var_id;
112 const auto uniq_name = ss.str();
113 SharedVar var{ var_link.group, uniq_name, runtime_desc };
114
115 if(var_link.group == SharedVarGroup::Argument)
116 {
117 _arguments.emplace(var_id, var);
118 _num_var++;
119 _var_id_lut[var_link.arg_id] = var_id;
120 }
121 else if(var_link.group == SharedVarGroup::Automatic)
122 {
123 if(var_link.io == SharedVarIO::Output)
124 {
125 _global_vars.emplace(var_id, var);
126 _num_var++;
127 _var_id_lut[var_link.arg_id] = var_id;
128 }
129 else
130 {
131 // For the input link, the var (and thus its arg_id) will always have been added by the time we get here if we traverse components in topological order
132 var = get_var(var_link.arg_id);
133 }
134 }
135 else
136 {
137 ARM_COMPUTE_ERROR("Unrecognised SharedVarGroup");
138 }
139 return var;
140 }
141
142 SharedVar get_var(ArgumentID arg_id) const
143 {
144 const auto var_id = _var_id_lut.at(arg_id); // arg_id has to exist in lut to begin with
145 auto it = _global_vars.find(var_id);
146 if(it != _global_vars.end())
147 {
148 return it->second;
149 }
150 it = _arguments.find(var_id);
151 if(it != _arguments.end())
152 {
153 return it->second;
154 }
155 ARM_COMPUTE_ERROR("Cannot find component variable");
156 }
157
158 /** @note The arguments are returned in the order they are added
159 */
160 Arguments get_kernel_arguments() const
161 {
162 Arguments args{};
163 for(const auto &a : _arguments)
164 {
165 args.push_back(a.second);
166 }
167 return args;
168 }
169
170private:
171 using VarID = int32_t;
172
173private:
174 std::map<VarID, SharedVar> _global_vars{};
175 std::map<VarID, SharedVar> _arguments{};
176 std::unordered_map<ArgumentID, VarID> _var_id_lut{};
177 VarID _num_var{ 0 };
178};
179
180enum class ComponentType
181{
182 Simple,
183 Complex,
184 Store
185};
186
187using ComponentID = int32_t;
188using ComponentList = std::vector<ComponentID>;
189class IClKernelComponent
190{
191public:
192 using Link = SharedVarLink;
193 using Tag = std::string;
194 struct TagVal
195 {
196 TagVal() = default;
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000197 TagVal(const SharedVarTable::SharedVar &var)
Giorgio Arena232c4522022-03-03 10:09:01 +0000198 : value{ var.uniq_name }
199 {
200 }
201
Gunes Bayir16c56972022-03-28 21:32:33 +0100202 template <typename T, ARM_COMPUTE_REQUIRES_TA(std::is_integral<T>::value)>
203 TagVal(T val)
204 : value{ support::cpp11::to_string(val) }
Giorgio Arena232c4522022-03-03 10:09:01 +0000205 {
206 }
207
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000208 TagVal(const std::string &val)
209 : value{ val }
210 {
211 }
212
Gunes Bayir16c56972022-03-28 21:32:33 +0100213 TagVal(const char *val)
214 : value{ std::string(val) }
215 {
216 }
217
218 TagVal(const DataType &data_type)
219 : value{ get_cl_type_from_data_type(data_type) }
220 {
221 }
222
Giorgio Arena232c4522022-03-03 10:09:01 +0000223 std::string value{};
224 };
225 using TagLUT = std::unordered_map<Tag, TagVal>; // Used to instantiating a code template / replacing tags
226public:
Gunes Bayir8a879832022-03-10 21:21:01 +0000227 IClKernelComponent(const ClKernelBlueprint *blueprint)
228 : _blueprint(blueprint)
229 {
230 }
231
232 ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE(IClKernelComponent);
233
Giorgio Arena232c4522022-03-03 10:09:01 +0000234 virtual ~IClKernelComponent() = default;
235 virtual ComponentType get_component_type() const = 0;
236 virtual std::vector<Link> get_links() const = 0;
237 virtual std::string name() const = 0;
238
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000239 // @note: some tags can be unused since they could be used only for the macros, or only for the component code
Giorgio Arena232c4522022-03-03 10:09:01 +0000240 static std::string replace_tags(const std::string &code_template, const TagLUT &tags)
241 {
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000242 std::string replaced_code = "";
243 bool scanning_pattern = false;
244 std::string pattern_found = "";
Giorgio Arena232c4522022-03-03 10:09:01 +0000245 for(size_t i = 0; i < code_template.size() - 1; ++i)
246 {
247 if(!scanning_pattern)
248 {
249 if(code_template[i] == '{' && code_template[i + 1] == '{')
250 {
251 i += 1;
252 scanning_pattern = true;
253 pattern_found = "";
254 }
255 else
256 {
257 replaced_code += code_template[i];
258 }
259 }
260 else
261 {
262 if(code_template[i] == '}' && code_template[i + 1] == '}')
263 {
264 i += 1;
265 scanning_pattern = false;
266 std::string err = "Pattern " + pattern_found + " not found in tags";
267 ARM_COMPUTE_ERROR_ON_MSG(tags.find(pattern_found) == tags.end(), err.c_str());
268 replaced_code += tags.find(pattern_found)->second.value;
Giorgio Arena232c4522022-03-03 10:09:01 +0000269 }
270 else
271 {
272 pattern_found += code_template[i];
273 }
274 }
275 }
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000276
Giorgio Arena232c4522022-03-03 10:09:01 +0000277 return replaced_code;
278 }
279 ComponentID id() const
280 {
281 return _id;
282 }
283 void set_id(ComponentID id)
284 {
285 _id = id;
286 }
287
288 virtual std::set<std::string> get_headers_list() const
289 {
290 return std::set<std::string> {};
291 }
292
293 virtual std::string get_additional_macros() const
294 {
295 return "";
296 }
297
298 virtual std::string get_component_code() const
299 {
300 return "";
301 }
Gunes Bayir8a879832022-03-10 21:21:01 +0000302
303 virtual Window get_window() const
304 {
305 return Window{};
306 }
Giorgio Arena232c4522022-03-03 10:09:01 +0000307 /** "Allocate" all shared variables used in a component to the @p vtable, and generate a TagLUT used to instantiate the component code
308 *
309 * @param vtable
310 * @return TagLUT
311 */
312 virtual TagLUT allocate_vars(SharedVarTable &vtable) const = 0;
313
314 virtual std::string get_dst_addr_calculation() const
315 {
316 return "";
317 }
318
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000319 virtual CLBuildOptions generate_build_options() const
320 {
321 return CLBuildOptions{};
322 }
323
Gunes Bayir8a879832022-03-10 21:21:01 +0000324protected:
325 const ClKernelBlueprint *_blueprint;
326
Giorgio Arena232c4522022-03-03 10:09:01 +0000327private:
328 ComponentID _id{};
329};
330
331using ComponentUniquePtr = std::unique_ptr<IClKernelComponent>;
332
333/** Intermediate representation of the final, complete kernel source.
334 */
335struct ClKernelBlueprint::Implementation
336{
337public:
338 Implementation() = default;
339 ~Implementation() = default;
340
341public:
342 ArgumentID add_kernel_argument(const ClTensorDescriptor &tensor_desc)
343 {
344 _kernel_arguments.insert(std::make_pair(_num_args, tensor_desc));
345 _shared_var_group_lut[_num_args] = SharedVarGroup::Argument;
346 return _num_args++;
347 }
348
349 ArgumentID add_intermediate_tensor()
350 {
351 _intermediate_tensors.insert(_num_args);
352 _shared_var_group_lut[_num_args] = SharedVarGroup::Automatic;
353 return _num_args++;
354 }
355
356 void set_tile_info(const TileDescriptor &tile_info)
357 {
358 _tile_info = tile_info;
359 }
360
361 SharedVarGroup group(ArgumentID arg_id) const
362 {
363 if(arg_id == g_arg_placeholder)
364 {
365 // In case of placeholder, don't care what we return;
366 return SharedVarGroup::Argument;
367 }
368 return _shared_var_group_lut.at(arg_id);
369 }
370
371 void validate_arg_ids(std::initializer_list<ArgumentID> args) const
372 {
373 for(const auto arg_id : args)
374 {
375 ARM_COMPUTE_UNUSED(arg_id);
376 ARM_COMPUTE_ERROR_ON_MSG(_kernel_arguments.find(arg_id) == _kernel_arguments.end() && _intermediate_tensors.find(arg_id) == _intermediate_tensors.end() && arg_id != g_arg_placeholder,
377 "Trying to use an argument that hasn't been added to the blueprint");
378 }
379 }
380
381 void add_component(ComponentUniquePtr component)
382 {
383 if(component->get_component_type() == ComponentType::Complex)
384 {
385 ++_num_complex_components;
386 ARM_COMPUTE_ERROR_ON_MSG(_num_complex_components > 1, "Only one complex component per blueprint is supported.");
387 }
388
389 // This flag specifies if the current component is the root of the component graph
390 // If the root is set to -1, it means that a root hasn't been added yet
391 bool is_graph_root = true;
392
393 // Get an unique ID for the component that's being added
394 const ComponentID component_id = _num_components++;
395 component->set_id(component_id);
396
397 // Add this component to the component graph. Don't connect it to anything yet
398 _component_graph.emplace(component_id, ComponentList{});
399
400 int32_t positional_arg = 0;
401
402 // For every { arg_id, arg_io } passed along with this component...
403 for(const auto &link : component->get_links())
404 {
405 const ArgumentID &arg_id = link.arg_id;
406 const SharedVarIO &arg_io = link.io;
407
408 // A component is considered root only if all its input arguments are kernel arguments (or placeholders, which means nullptr)
409 // This performs a check on every argument, and if one of them doesn't respect the condition, the component is not considered root
410 is_graph_root &= (_kernel_arguments.find(arg_id) != _kernel_arguments.end()) || (arg_io == SharedVarIO::Output) || (arg_id == g_arg_placeholder);
411
412 // Add the arg_id to the map describing the input/output relationship between an argument and the components that use it, if it doesn't yet exist there
413 if(_outgoing_components.find(arg_id) == _outgoing_components.end())
414 {
415 _outgoing_components.emplace(arg_id, ComponentList{});
416 _incoming_components.emplace(arg_id, ComponentList{});
417 }
418
419 // If it's an input argument, connect any other component that has it as output with this component
420 // Additionally, set this component as one that treats this argument as "Input" (append to index 0)
421 // This is used so that we keep track of whether two components use the same argument, one as input and one as output
422 if(arg_io == SharedVarIO::Input)
423 {
424 for(const auto &prev_component : _incoming_components[arg_id])
425 {
426 _component_graph[prev_component].push_back(component_id);
427 }
428
429 _outgoing_components[arg_id].push_back(component_id);
430 }
431 // If it's an output argument, connect this component with any other component that has it as input
432 // Additionally, set this component as one that treats this argument as "Output" (append to index 1)
433 else
434 {
Gunes Bayir8a879832022-03-10 21:21:01 +0000435 if(component->get_component_type() == ComponentType::Store)
436 {
437 ARM_COMPUTE_ERROR_ON_MSG(_dst_id >= 0, "Trying to add more than one dst argument to the graph");
438 _dst_id = arg_id;
439 }
440
Giorgio Arena232c4522022-03-03 10:09:01 +0000441 for(const auto &subseq_component : _outgoing_components[arg_id])
442 {
443 _component_graph[component_id].push_back(subseq_component);
444 }
445
446 _incoming_components[arg_id].push_back(component_id);
447 }
448
449 ++positional_arg;
450 }
451
452 if(is_graph_root)
453 {
454 ARM_COMPUTE_ERROR_ON_MSG(_graph_root >= 0, "Trying to add more than one root to the graph");
455 _graph_root = component_id;
456 }
457
458 // Finally, add this component to the dictionary of components
459 _components.insert(std::make_pair(component_id, std::move(component)));
460 }
461
462 std::string build_kernel_name() const
463 {
464 std::string name = "";
465
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000466 traverse([&](std::stack<ComponentID> stack)
Giorgio Arena232c4522022-03-03 10:09:01 +0000467 {
468 name += _components.find(stack.top())->second->name() + (stack.size() > 2 ? "___" : "");
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000469 });
Giorgio Arena232c4522022-03-03 10:09:01 +0000470
Giorgio Arena232c4522022-03-03 10:09:01 +0000471 return name;
472 }
473
474 std::string build_code()
475 {
476 ARM_COMPUTE_ERROR_ON_MSG(_graph_root < 0, "No root found in the component graph");
477
478 // These data structures will hold the data from all the components in the blueprint
479 std::set<std::string> headers_list{};
480 std::set<std::string> additional_macros{};
481 std::vector<std::string> component_codes{}; // vector because order matters
482
483 // Go through the components graph (topological sort) and fill the data structures above
484 auto stack = topological_sort();
485 while(!stack.empty())
486 {
487 auto curr_component_id = stack.top();
488 auto &curr_component = _components.find(curr_component_id)->second;
489
490 auto curr_headers_list = curr_component->get_headers_list();
491 auto curr_additional_macros = curr_component->get_additional_macros();
492 auto curr_component_code = curr_component->get_component_code();
493 const auto var_lut = curr_component->allocate_vars(_vtable); // Ideally can be merged with get_component_code once we have finer-grained code generation technique
494 component_codes.push_back(IClKernelComponent::replace_tags(curr_component_code, var_lut));
495
496 headers_list.insert(curr_headers_list.begin(), curr_headers_list.end());
497 if(!curr_additional_macros.empty()) // Some components might not have any
498 {
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000499 additional_macros.insert(IClKernelComponent::replace_tags(curr_additional_macros, var_lut));
Giorgio Arena232c4522022-03-03 10:09:01 +0000500 }
501
502 stack.pop();
503 }
504
505 // This section assembles the data gathered by traversing the graph into the string "code"
506 std::string code = "";
507
508 for(auto &header : headers_list)
509 {
Giorgio Arena892b70a2022-03-30 12:23:10 +0100510#if defined(EMBEDDED_KERNELS)
511 code += CLKernelLibrary::get().get_program(header).first;
512#else // defined(EMBEDDED_KERNELS)
Giorgio Arena232c4522022-03-03 10:09:01 +0000513 code += "#include \"" + header + "\"\n";
Giorgio Arena892b70a2022-03-30 12:23:10 +0100514#endif // defined(EMBEDDED_KERNELS)
Giorgio Arena232c4522022-03-03 10:09:01 +0000515 }
516
517 for(auto &macros : additional_macros)
518 {
519 code += macros;
520 }
521
522 code += generate_kernel_signature(_vtable.get_kernel_arguments());
523
524 code += "\n{\n\n";
525
526 code += " //------------------ START KERNEL_BUILDER_COORDINATE ---------------------\n\n";
527 code += generate_global_section();
528 code += " //------------------ END KERNEL_BUILDER_COORDINATE ---------------------\n";
529
530 for(auto &component_code : component_codes)
531 {
532 code += component_code;
533 }
534
535 code += "}\n";
536
537 return code;
538 }
539
540 std::string build_config_id() const
541 {
542 return "";
543 }
544
545 CLBuildOptions build_options() const
546 {
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000547 CLBuildOptions build_opts{};
548
549 traverse([&](std::stack<ComponentID> stack)
550 {
551 build_opts.add_options(_components.find(stack.top())->second->generate_build_options().options());
552 });
553
554 return build_opts;
555 }
556
557 TileDescriptor get_tile_info() const
558 {
559 return _tile_info;
Giorgio Arena232c4522022-03-03 10:09:01 +0000560 }
561
562 Window get_execution_window() const
563 {
Gunes Bayir8a879832022-03-10 21:21:01 +0000564 ARM_COMPUTE_ERROR_ON_MSG(_graph_root < 0, "No root found in the component graph");
565 ARM_COMPUTE_ERROR_ON_MSG(_dst_id == -1, "Destination Tensor Id should be ready before calling get_execution_window()");
566
567 return _components.find(_graph_root)->second->get_window();
568 }
569
570 ArgumentID get_dst_id() const
571 {
572 return _dst_id;
Giorgio Arena232c4522022-03-03 10:09:01 +0000573 }
574
575 ClKernelArgList get_arguments() const
576 {
577 ClKernelArgList arg_list{};
578 for(const auto &arg_var : _vtable.get_kernel_arguments())
579 {
580 arg_list.push_back(arg_var.desc);
581 }
582 return arg_list;
583 }
584
Gunes Bayir8a879832022-03-10 21:21:01 +0000585 const ClTensorDescriptor *get_kernel_argument(const ArgumentID id) const
586 {
587 auto it = _kernel_arguments.find(id);
588 if(it != _kernel_arguments.end())
589 {
590 return &_kernel_arguments.find(id)->second;
591 }
592 return nullptr;
593 }
594
595 ITensorInfo *get_kernel_argument_info(const ArgumentID id) const
596 {
597 const ClTensorDescriptor *arg_desc = get_kernel_argument(id);
598 if(arg_desc != nullptr)
599 {
600 return arg_desc->tensor_info;
601 }
602 return nullptr;
603 }
604
Giorgio Arena232c4522022-03-03 10:09:01 +0000605private:
606 void topological_sort_utility(ComponentID component_id, std::unordered_set<ComponentID> &visited, std::stack<ComponentID> &stack) const
607 {
608 visited.insert(component_id);
609
610 for(auto connected_component : _component_graph.find(component_id)->second)
611 {
612 if(visited.find(connected_component) == visited.end())
613 {
614 topological_sort_utility(connected_component, visited, stack);
615 }
616 }
617
618 stack.push(component_id);
619 }
620
621 std::stack<ComponentID> topological_sort() const
622 {
623 std::stack<ComponentID> stack{};
624 std::unordered_set<ComponentID> visited{};
625
626 topological_sort_utility(_graph_root, visited, stack);
627
628 return stack;
629 }
630
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000631 void traverse(const std::function<void(std::stack<ComponentID>)> &func) const
632 {
633 std::stack<ComponentID> stack = topological_sort();
634
635 while(!stack.empty())
636 {
637 func(stack);
638 stack.pop();
639 }
640 }
641
Giorgio Arena232c4522022-03-03 10:09:01 +0000642 std::string generate_argument_declaration(const SharedVarTable::SharedVar &var) const
643 {
644 ARM_COMPUTE_ERROR_ON_MSG(var.group != SharedVarGroup::Argument, "An argument declaration can only be generated from a kernel argument");
645 std::string code;
646 switch(var.desc.tensor_arg_type)
647 {
Gunes Bayir16c56972022-03-28 21:32:33 +0100648 case TensorArgType::Vector:
649 {
650 code += "\n VECTOR_DECLARATION(" + var.uniq_name + ")";
651 break;
652 }
Giorgio Arena232c4522022-03-03 10:09:01 +0000653 case TensorArgType::Image:
654 {
Gunes Bayir16c56972022-03-28 21:32:33 +0100655 code += "\n IMAGE_DECLARATION(" + var.uniq_name + ")";
Giorgio Arena232c4522022-03-03 10:09:01 +0000656 break;
657 }
658 case TensorArgType::Image_3D:
659 {
Gunes Bayir16c56972022-03-28 21:32:33 +0100660 code += "\n IMAGE_DECLARATION(" + var.uniq_name + "),";
661 code += "\n uint " + var.uniq_name + "_stride_z";
Giorgio Arena232c4522022-03-03 10:09:01 +0000662 break;
663 }
664 case TensorArgType::Image_3D_Export_To_ClImage2D:
665 {
Gunes Bayir16c56972022-03-28 21:32:33 +0100666 code += "\n __read_only image2d_t " + var.uniq_name + "_img,";
667 code += "\n uint " + var.uniq_name + "_stride_z";
668 break;
669 }
670 case TensorArgType::Tensor_4D_t_Buffer:
671 {
672 code += "\n TENSOR4D_T(" + var.uniq_name + ", BUFFER)";
673 break;
674 }
675 case TensorArgType::Tensor_4D_t_Image:
676 {
677 code += "\n TENSOR4D_T(" + var.uniq_name + ", IMAGE)";
Giorgio Arena232c4522022-03-03 10:09:01 +0000678 break;
679 }
680 default:
681 {
682 ARM_COMPUTE_ERROR("Unsupported declaration generation for TensorArgType");
683 }
684 }
685 return code;
686 }
687
688 std::string generate_kernel_signature(const SharedVarTable::Arguments &argument_list) const
689 {
690 std::string code = "\n__kernel void " + build_kernel_name() + "(";
691
692 for(const auto &arg : argument_list)
693 {
Gunes Bayir16c56972022-03-28 21:32:33 +0100694 code += generate_argument_declaration(arg) + ",";
Giorgio Arena232c4522022-03-03 10:09:01 +0000695 }
696
697 code[code.length() - 1] = ')';
698
699 return code;
700 }
701
702 std::string generate_global_section() const
703 {
Gunes Bayir16c56972022-03-28 21:32:33 +0100704 std::string code = "";
705 code += " uint g_x = get_global_id(0);\n";
Giorgio Arena232c4522022-03-03 10:09:01 +0000706 code += " uint g_y = get_global_id(1);\n";
707 code += " uint g_z = get_global_id(2);\n\n";
708
709 size_t tile_dim_x = _tile_info.empty() ? 1 : _tile_info.tile_dims.x();
710 size_t tile_dim_y = _tile_info.empty() ? 1 : _tile_info.tile_dims.y();
711
712 switch(_tile_info.clipping)
713 {
714 case ClippingStrategy::TOP_LEFT:
715 code += " const bool g_cond_x = (g_x == 0);\n";
716 code += " const bool g_cond_y = (g_y == 0);\n";
717 break;
718 case ClippingStrategy::TOP_RIGHT:
719 code += " const bool g_cond_x = ((g_x + 1) * " + std::to_string(tile_dim_x) + " >= " + std::to_string(_tile_info.boundaries.x()) + ");\n";
720 code += " const bool g_cond_y = (g_y == 0);\n";
721 break;
722 case ClippingStrategy::BOTTOM_LEFT:
723 code += " const bool g_cond_x = (g_x == 0);\n";
724 code += " const bool g_cond_y = ((g_y + 1) * " + std::to_string(tile_dim_y) + " >= " + std::to_string(_tile_info.boundaries.y()) + ");\n";
725 break;
726 case ClippingStrategy::BOTTOM_RIGHT:
727 code += " const bool g_cond_x = ((g_x + 1) * " + std::to_string(tile_dim_x) + " >= " + std::to_string(_tile_info.boundaries.x()) + ");\n";
728 code += " const bool g_cond_y = ((g_y + 1) * " + std::to_string(tile_dim_y) + " >= " + std::to_string(_tile_info.boundaries.y()) + ");\n";
729 break;
730 default:
731 ARM_COMPUTE_ERROR("Unsupported clipping strategy");
732 }
733
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000734 code += "\n REPEAT_VAR_INIT_TO_CONST(" + std::to_string(tile_dim_y) + ", uint, g_zout, 0);\n";
Giorgio Arena232c4522022-03-03 10:09:01 +0000735 code += " REPEAT_VAR_INIT_TO_CONST(16, uint, g_zero, 0);\n\n";
736
737 return code;
738 }
739
740 TileDescriptor _tile_info{};
741
742 int32_t _num_args{};
743 int32_t _num_components{};
744 int32_t _num_complex_components{};
745
Giorgio Arenabd44caa2022-03-15 13:45:15 +0000746 ArgumentID _dst_id{ -1 }; // Initially set to -1, which means the graph has no dst yet, since node IDs are positive numbers
Gunes Bayir8a879832022-03-10 21:21:01 +0000747
Giorgio Arena232c4522022-03-03 10:09:01 +0000748 // Argument, components and intermediate tensors IDs with corresponding ptrs (except intermediate)
749 std::unordered_map<ComponentID, ComponentUniquePtr> _components{};
750 std::unordered_map<ArgumentID, ClTensorDescriptor> _kernel_arguments{};
751 std::unordered_set<ArgumentID> _intermediate_tensors{};
752 // Argument group lookup. Can be replaced by extending the ArgumentID type to include group info
753 std::unordered_map<ArgumentID, SharedVarGroup> _shared_var_group_lut{};
754
755 // Tracks all variables (e.g.: kernel arguments, kernel "global variables")
756 SharedVarTable _vtable{};
757
758 // Component directed graph (represented by an adjecency list of Component IDs)
759 // This is used to understand the ordering and bindings between components when generating the kernel
760 // It's initially set to -1 which means the graph has no root yet, since node IDs are positive numbers
761 ComponentID _graph_root{ -1 };
762 std::unordered_map<ComponentID, ComponentList> _component_graph{};
763
764 // Additional data structures used to define the relationships between components and arguments
765 // For each argument, it contains the list of components that consider it as an incoming or an outgoing argument
766 // E.g. tensor0 -> component0 -> tensor1
767 // _outgoing_components[tensor0] == {component0} (component0 is the outgoing component of tensor0. Component0 treats tensor0 as an input tensor)
768 // _incoming_components[tensor1] == {component0} (component0 is the incoming component of tensor1. Component1 treats tensor1 as an output tensor)
769 std::unordered_map<ArgumentID, ComponentList> _outgoing_components{};
770 std::unordered_map<ArgumentID, ComponentList> _incoming_components{};
771};
772
773} // namespace dynamic_fusion
774} // namespace experimental
775} // namespace arm_compute
776#endif //ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMMON_H
777
778#endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)