1 files changed, 59 insertions, 2 deletions
diff --git a/compute_kernel_writer/src/cl/CLKernelWriter.cpp b/compute_kernel_writer/src/cl/CLKernelWriter.cpp
index 9363076901..88ada37d71 100644
--- a/compute_kernel_writer/src/cl/CLKernelWriter.cpp
+++ b/compute_kernel_writer/src/cl/CLKernelWriter.cpp
@@ -26,6 +26,8 @@
 #include "ckw/Error.h"
 #include "ckw/Kernel.h"
 #include "ckw/TileOperand.h"
+#include "ckw/types/TargetLanguage.h"
+#include "src/ITensorComponent.h"
 #include "src/cl/CLHelpers.h"
 #include "src/cl/CLTensorArgument.h"
 #include "src/cl/CLTile.h"
@@ -39,8 +41,63 @@ CLKernelWriter::~CLKernelWriter() = default;
 
 std::unique_ptr<Kernel> CLKernelWriter::emit_kernel(const std::string &name)
 {
-    CKW_UNUSED(name);
-    CKW_THROW_MSG("Not implemented!");
+    std::string code;
+
+    code += "__kernel void ";
+    code += name;
+    code += "\n(\n";
+
+    // Create the list of arguments.
+    std::vector<KernelArgument> arguments;
+
+    for(const auto &tensor : _tensors)
+    {
+        const auto tensor_id = tensor->info().id();
+
+        const auto storages   = tensor->storages();
+        const auto components = tensor->components();
+
+        for(const auto &storage : storages)
+        {
+            code += cl_get_variable_storagetype_as_string(storage.type);
+            code += " ";
+            code += storage.val;
+            code += ",\n";
+
+            arguments.emplace_back(tensor_id, storage.type);
+        }
+
+        for(const auto &component : components)
+        {
+            const auto &tile      = component->tile();
+            const auto &tile_info = tile.info();
+
+            CKW_ASSERT(tile_info.height() == 1);
+            CKW_ASSERT(tile_info.width() == 1);
+
+            code += cl_get_variable_datatype_as_string(tile_info.data_type(), 1);
+            code += " ";
+            code += tile.name();
+            code += ",\n";
+
+            arguments.emplace_back(tensor_id, component->component_type());
+        }
+    }
+
+    if(code.size() >= 2 && code[code.size() - 2] == ',' && code[code.size() - 1] == '\n')
+    {
+        // Remove the last comma in the argument list.
+        code.pop_back();
+        code[code.size() - 1] = '\n';
+    }
+
+    code += ")\n{\n";
+
+    code += _body_source_code;
+
+    code += "}\n";
+
+    return std::make_unique<Kernel>(TargetLanguage::OpenCL, arguments, code);
 }
 
 void CLKernelWriter::comment(const std::string &text)