Merge pull request #3 from spcl/python_codegen

Python codegen

Author: TizianoDeMatteis

README.md

@@ -14,25 +14,21 @@ For more details, see our [paper](https://arxiv.org/abs/1907.07929).
 
 The library depends on:
 
-* Intel FPGA SDK for OpenCL pro, version 18+ ([http://fpgasoftware.intel.com/opencl/](http://fpgasoftware.intel.com/opencl/))
+* Intel FPGA SDK for OpenCL pro, version 19+ ([http://fpgasoftware.intel.com/opencl/](http://fpgasoftware.intel.com/opencl/))
 * GCC (version 5+)
 * Rapidjson ([http://rapidjson.org/](http://rapidjson.org/))
 * Google Test (only for unit tests)
+* Python 3.6+
 
 ### Installation
 
 After cloning this repository, make sure you clone the [rapidjson](http://rapidjson.org/) submodule dependency, by executing the following command:
 
 ```
 git submodule update --init
+pip install -r codegen/requirements.txt
 ```
 
-After this, the included Makefile can be used to compile code and modules generator:
-
-
-```
-make all
-```
 
 ## The FBLAS library
 
@@ -47,19 +43,25 @@ the device, she can invoke the desired FBLAS routines working on the FPGA memory
 
 For further information on how to use the library, please refer to the [wiki](https://github.com/spcl/FBLAS/wiki).
 
-
 ## Publication
 If you use FBLAS, please cite us:
 ```
-@article{
-  author={Tiziano De Matteis and Johannes de Fine Licht and Torsten Hoefler},
-  title={{FBLAS: Streaming Linear Algebra on FPGA}},
-  journal={CoRR},
-  year={2019},
-  month={Jul.},
-  volume={abs/1907.07929},
+
+@inproceedings{fblas,
+    author = {De Matteis, Tiziano and de Fine Licht, Johannes and Hoefler, Torsten},
+    title = {FBLAS: Streaming Linear Algebra on FPGA},
+    year = {2020},
+    isbn = {9781728199986},
+    publisher = {IEEE Press},
+    booktitle = {Proceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis},
+    articleno = {59},
+    numpages = {13},
+    keywords = {high level synthesis, spatial architectures, hardware library},
+    location = {Atlanta, Georgia},
+    series = {SC '20}
 }
 ```
+Concerning the Artifact Evaluation of the paper, you will find detailed information in subfolder `evaluation`.
 
 
 ## Contact

codegen/.my_routine.cl

@@ -0,0 +1,82 @@
+ /**
+    FBLAS: BLAS implementation for Intel FPGA
+    Copyright (c) 2020 ETH-Zurich. All rights reserved.
+    See LICENSE for license information.
+
+    DOT performs the dot product of two vectors.
+
+    Streamed version: data is received from two input streams
+    CHANNEL_VECTOR_X and CHANNEL_VECTOR_Y having the proper type.
+    Data elements must be streamed with a padding equal to W
+    (padding data must be set to zero).
+
+    Result is streamed in an output channel at the end of the computation
+    in a channel CHANNEL_OUT
+
+*/
+
+#pragma OPENCL EXTENSION cl_intel_channels : enable
+
+
+#define W 64		    //width: number of multiplications performed per clock cycle
+
+//namings
+#define my_routine streaming_dot
+
+//channels names
+#define CHANNEL_VECTOR_X channel_gen_x
+#define CHANNEL_VECTOR_Y channel_gen_y
+#define CHANNEL_OUT channel_sink
+
+//architecture
+#define __STRATIX_10__
+
+//FBLAS_PARAMETERS_END
+
+#include <commons.h>
+
+channel float CHANNEL_VECTOR_X __attribute__((depth(1)));
+channel float CHANNEL_VECTOR_Y __attribute__((depth(1)));
+channel float CHANNEL_OUT __attribute__((depth(1)));
+
+
+/**
+    Performs streaming dot product: data is received through
+    CHANNEL_VECTOR_X and CHANNEL_VECTOR_Y. Result is sent
+    to CHANNEL_OUT.
+*/
+__kernel void my_routine(int N)
+{
+    __constant uint WIDTH = 1
+    float acc_o=0;
+    if(N>0)
+    {
+
+        const int outer_loop_limit=1+(int)((N-1)/WIDTH); //ceiling
+        float x[WIDTH],y[WIDTH];
+
+
+        //Strip mine the computation loop to exploit unrolling
+        for(int i=0; i<outer_loop_limit; i++)
+        {
+
+            float acc_i=0;
+            #pragma unroll
+            for(int j=0;j<WIDTH;j++)
+            {
+                x[j]=read_channel_intel(CHANNEL_VECTOR_X);
+                y[j]=read_channel_intel(CHANNEL_VECTOR_Y);
+                acc_i+=x[j]*y[j];
+
+            }
+
+                acc_o+=acc_i;
+
+        }
+
+    }
+    else //no computation: result is zero
+        acc_o=0.0f;
+    //write to the sink
+    write_channel_intel(CHANNEL_OUT,acc_o);
+}

codegen/codegen/__init__.py

@@ -0,0 +1,46 @@
+{
+    "helper": [{
+            "name": "read_vector_x",
+            "optional_parameters": ["width", "stride"]
+        },
+        {
+            "name": "read_vector_y",
+            "optional_parameters": ["width", "stride"]
+        },
+        {
+            "name": "generate_dummy_vector",
+            "optional_parameters": ["width"]
+        },
+        {
+            "name": "write_vector",
+            "optional_parameters": ["width", "stride"]
+        },
+        {
+            "name": "vector_sink",
+            "optional_parameters": ["width"]
+        },
+        {
+            "name": "write_scalar"
+        },
+        {
+            "name": "read_matrix",
+            "required_parameters": ["tiles order", "elements order"],
+            "optional_parameters": ["tile N size", "tile M size", "width"]
+        },
+        {
+            "name": "write_matrix",
+            "required_parameters": ["tiles order", "elements order"],
+            "optional_parameters": ["tile N size", "tile M size", "width"]
+        },
+        {
+            "name": "read_matrix_4_modules",
+            "required_parameters": ["tiles order", "elements order"],
+            "optional_parameters": ["tile N size", "tile M size", "width"]
+        },
+        {
+            "name": "generate_dummy_matrix",
+            "required_parameters": ["tiles order", "elements order"],
+            "optional_parameters": ["tile N size", "tile M size", "width"]
+        }
+    ]
+}

codegen/codegen/conf/helper_defs_modules_codegen.json

@@ -88,17 +88,17 @@
         {
             "name": "gemm",
             "required_parameters": ["order", "transa", "transb"],
-            "optional_parameters": ["computational width x", "computational width y" , "tile size", "systolic"]
+            "optional_parameters": ["width x", "width y" , "tile size", "systolic", "vect size"]
         },
 		{
 			"name": "syrk",
 			"required_parameters": ["order", "trans", "uplo"],
-			"optional_parameters": ["computational width x", "computational width y" , "tile size"]
+			"optional_parameters": ["width x", "width y" , "tile size"]
 		},
         {
             "name": "syr2k",
             "required_parameters": ["order", "trans", "uplo"],
-            "optional_parameters": ["computational width x", "computational width y" , "tile size"]
+            "optional_parameters": ["width x", "width y" , "tile size"]
         },
         {
             "name": "trsm",

…nerator/routines_definitions_tier_2.json …/codegen/conf/routine_defs_host_api.jsoninclude/generator/routines_definitions_tier_2.json renamed to codegen/codegen/conf/routine_defs_host_api.json include/generator/routines_definitions_tier_2.json renamed to codegen/codegen/conf/routine_defs_host_api.json

@@ -2,27 +2,28 @@
     "routine": [{
             "name": "dot",
             "optional_parameters": ["width"],
-            "required_inputs": ["x", "y"],
-            "required_outputs": ["res"]
+            "required_inputs": ["in_x", "in_y"],
+            "required_outputs": ["out_res"]
         },
         {
             "name": "axpy",
             "optional_parameters": ["width"],
-            "required_inputs": ["x", "y"],
-            "required_outputs": ["res"]
+            "required_inputs": ["in_x", "in_y"],
+            "required_outputs": ["out_res"]
         },
         {
             "name": "gemv",
             "required_parameters": ["trans"],
             "optional_parameters": ["width", "tile N size", "tile M size", "A tiles order", "A elements order"],
-            "required_inputs": ["x", "y","A"],
-            "required_outputs": ["res"]
+            "required_inputs": ["in_x", "in_y","in_A"],
+            "required_outputs": ["out_res"],
+            "optional_outputs": ["out_y_updates"]
         },
         {
             "name": "ger",
             "optional_parameters": ["width", "tile N size", "tile M size", "A tiles order", "A elements order"],
-            "required_inputs": ["x", "y","A"],
-            "required_outputs": ["res"]
+            "required_inputs": ["in_x", "in_y","in_A"],
+            "required_outputs": ["out_res"]
         }
     ]
 }