1..  Licensed to the Apache Software Foundation (ASF) under one
2    or more contributor license agreements.  See the NOTICE file
3    distributed with this work for additional information
4    regarding copyright ownership.  The ASF licenses this file
5    to you under the Apache License, Version 2.0 (the
6    "License"); you may not use this file except in compliance
7    with the License.  You may obtain a copy of the License at
8
9..    http://www.apache.org/licenses/LICENSE-2.0
10
11..  Unless required by applicable law or agreed to in writing,
12    software distributed under the License is distributed on an
13    "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
14    KIND, either express or implied.  See the License for the
15    specific language governing permissions and limitations
16    under the License.
17
18
19HLS Backend Example
20===================
21
22TVM supports Xilinx FPGA board with SDAccel.  Here is a tutorial for how to deploy TVM to AWS F1 FPGA instance.
23
24.. note::
25
26    This feature is still experimental.  We cannot use SDAccel to deploy an end to end neural networks for now.
27
28We use two python scripts for this tutorial.
29
30- build.py - a script to synthesize FPGA bitstream.
31
32  .. code:: python
33
34      import tvm
35      from tvm import te
36
37      tgt_host="llvm"
38      tgt="sdaccel"
39
40      n = te.var("n")
41      A = te.placeholder((n,), name='A')
42      B = te.placeholder((n,), name='B')
43      C = te.compute(A.shape, lambda i: A[i] + B[i], name="C")
44
45      s = te.create_schedule(C.op)
46      px, x = s[C].split(C.op.axis[0], nparts=1)
47
48      s[C].bind(px, tvm.te.thread_axis("pipeline"))
49
50      fadd = tvm.build(s, [A, B, C], tgt, target_host=tgt_host, name="myadd")
51
52      fadd.save("myadd.o")
53      fadd.imported_modules[0].save("myadd.xclbin")
54
55      tvm.contrib.cc.create_shared("myadd.so", ["myadd.o"])
56
57- run.py - a script to use FPGA as an accelerator.
58
59  .. code:: python
60
61      import tvm
62      import numpy as np
63      import os
64
65      tgt="sdaccel"
66
67      fadd = tvm.runtime.load_module("myadd.so")
68      if os.environ.get("XCL_EMULATION_MODE"):
69          fadd_dev = tvm.runtime.load_module("myadd.xclbin")
70      else:
71          fadd_dev = tvm.runtime.load_module("myadd.awsxclbin")
72      fadd.import_module(fadd_dev)
73
74      ctx = tvm.context(tgt, 0)
75
76      n = 1024
77      a = tvm.nd.array(np.random.uniform(size=n).astype("float32"), ctx)
78      b = tvm.nd.array(np.random.uniform(size=n).astype("float32"), ctx)
79      c = tvm.nd.array(np.zeros(n, dtype="float32"), ctx)
80
81      fadd(a, b, c)
82      tvm.testing.assert_allclose(c.asnumpy(), a.asnumpy() + b.asnumpy())
83
84
85Setup
86-----
87
88- Launch an instance using the FPGA Developer AMI.  We don't need an F1 instance for emulation and synthesis, so it is recommended to use a lower cost instance for them.
89- Setup AWS FPGA development kit.
90
91  .. code:: bash
92
93      git clone https://github.com/aws/aws-fpga.git
94      cd aws-fpga
95      source sdaccel_setup.sh
96      source ${XILINX_SDX}/settings64.sh
97
98- Setup TVM with OpenCL enabled.
99
100Emulation
101---------
102
103- Create emconfig.json for emulation.
104
105  .. code:: bash
106
107      emconfigutil --platform ${AWS_PLATFORM} --nd 1
108
109- Copy emconfig.json to the python binary directory.  It is because the current Xilinx toolkit assumes that both host binary and the emconfig.json file are in the same path.
110
111  .. code:: bash
112
113      cp emconfig.json $(dirname $(which python))
114
115- Run software emulation
116
117  .. code:: bash
118
119      export XCL_EMULATION_MODE=1
120      export XCL_TARGET=sw_emu
121
122      python build.py
123      python run.py
124
125- Run hardware emulation
126
127  .. code:: bash
128
129      export XCL_EMULATION_MODE=1
130      export XCL_TARGET=hw_emu
131
132      python build.py
133      python run.py
134
135Synthesis
136---------
137
138- Run synthesis with the following script.
139
140  .. code:: bash
141
142      unset XCL_EMULATION_MODE
143      export XCL_TARGET=hw
144
145      python build.py
146
147- Create AWS FPGA image and upload it to AWS S3.
148
149  .. code:: bash
150
151      ${SDACCEL_DIR}/tools/create_sdaccel_afi.sh \
152          -xclbin=myadd.xclbin -o=myadd \
153          -s3_bucket=<bucket-name> -s3_dcp_key=<dcp-folder-name> \
154          -s3_logs_key=<logs-folder-name>
155
156  This also generates an awsxclbin file, which is necessary to use the AWS FPGA image on F1 instances.
157
158Run
159---
160
161- Launch Amazon EC2 F1 instance.
162- Copy ``myadd.so``, ``myadd.awsxclbin``, and ``run.py`` to the F1 instance.
163- Setup AWS FPGA development kit.
164
165  .. code:: bash
166
167      git clone https://github.com/aws/aws-fpga.git
168      cd aws-fpga
169      source sdaccel_setup.sh
170
171- Setup TVM with OpenCL enabled.
172- Become root and setup environment variables.
173
174  .. code:: bash
175
176      sudo sh
177      source ${INSTALL_ROOT}/setup.sh
178
179- Run
180
181  .. code:: bash
182
183      python run.py
184