diff --git a/applications/ta2/designs/ta2_unb2b_mm_demo/Makefile b/applications/ta2/designs/ta2_unb2b_mm_demo/Makefile
new file mode 100644
index 0000000000000000000000000000000000000000..131d1b1dcc12ee06b35992e23d72bdd45136556f
--- /dev/null
+++ b/applications/ta2/designs/ta2_unb2b_mm_demo/Makefile
@@ -0,0 +1,120 @@
+######################
+### SETUP ###
+######################
+ifeq ($(VERBOSE),1)
+ECHO :=
+else
+ECHO := @
+endif
+
+# Where is the Intel(R) FPGA SDK for OpenCL(TM) software?
+ifeq ($(wildcard $(INTELFPGAOCLSDKROOT)),)
+$(error Set INTELFPGAOCLSDKROOT to the root directory of the Intel(R) FPGA SDK for OpenCL(TM) software installation)
+endif
+ifeq ($(wildcard $(INTELFPGAOCLSDKROOT)/host/include/CL/opencl.h),)
+$(error Set INTELFPGAOCLSDKROOT to the root directory of the Intel(R) FPGA SDK for OpenCL(TM) software installation.)
+endif
+
+###########################
+### Basic configuration ###
+###########################
+
+# Name of unb2b BSP
+UNB2B_BSP=ta2_unb2b_bsp
+
+# Compile directory
+BUILDDIR=$(RADIOHDL_BUILD_DIR)/unb2b/OpenCL/$(lastword $(subst /, ,$(abspath $(dir $(lastword $(MAKEFILE_LIST))))))
+
+
+##############################
+### Advanced Configuration ###
+##############################
+
+CXX= g++ #-mcmodel=medium
+CXXFLAGS= -std=c++11 -mavx2 -g -O3 -fopenmp #-DCL_ALTERA
+AOC= aoc
+AOCFLAGS= -v -g
+#AOCRFLAGS+= -fp-relaxed
+AOCRFLAGS+= -report
+AOCRFLAGS+= -opt-arg=-allow-io-channel-autorun-kernel
+#AOCRFLAGS+= -board=p385a_min_ax115_1710240
+AOCOFLAGS+= -board=$(UNB2B_BSP)
+
+AOCOFLAGS+= -I$(INTELOCLSDKROOT)/include/kernel_headers
+AOCXFLAGS+= -bsp-flow=flat
+ifneq ("$(SEED)", "")
+AOCXFLAGS+= -seed=$(SEED)
+endif
+INCLUDES= $(shell aocl compile-config) #-I..
+LDFLAGS= $(shell aocl link-config) #-ldl -lacl_emulator_kernel_rt #-lbfd
+CXXFLAGS+= $(INCLUDES)
+
+### Emulator configuration
+# Emulation Compilation flags
+ifeq ($(DEBUG),1)
+EMUCXXFLAGS += -g
+else
+EMUCXXFLAGS += -O2
+endif
+
+# Target
+TARGET := host
+TARGET_DIR := $(BUILDDIR)/bin
+
+# Directories
+INC_DIRS := host/lib/common/inc
+LIB_DIRS :=
+
+# Files
+INCS := $(wildcard )
+SRCS := $(wildcard host/src/*.cpp host/lib/common/src/*.cpp host/lib/common/src/AOCLUtils/*.cpp)
+LIBS := rt pthread
+
+### Emulator compilation
+# Make it all!
+%: %.cl $(TARGET_DIR)/$(TARGET)
+ (unset DISPLAY; mkdir -p $(BUILDDIR)/$* && $(AOC) -march=emulator -DEMULATOR $< -o $(TARGET_DIR)/$@.aocx -legacy-emulator $(AOCOFLAGS) $(AOCRFLAGS))
+
+# Host executable target.
+$(TARGET_DIR)/$(TARGET) : Makefile $(SRCS) $(INCS) $(TARGET_DIR)
+ $(ECHO)$(CXX) $(CPPFLAGS) $(CXXFLAGS) -fPIC $(foreach D,$(INC_DIRS),-I$D) \
+ $(INCLUDES) $(SRCS) $(LDFLAGS) \
+ $(foreach D,$(LIB_DIRS),-L$D) \
+ $(foreach L,$(LIBS),-l$L) \
+ -o $(TARGET_DIR)/$(TARGET)
+
+$(TARGET_DIR) :
+ $(ECHO)mkdir -p $(TARGET_DIR)
+
+# Standard make targets
+clean :
+ $(ECHO)rm -rf $(TARGET_DIR)/*
+
+
+### Device compilation
+%.d: %.cc
+ -$(CXX) $(CXXFLAGS) -MM -MT $@ -MT ${@:%.d=%.o} $< -o $@
+
+%.o: %.cc
+ $(CXX) -c $(CXXFLAGS) -o $@ $<
+
+%.aoco: %.cl
+ (unset DISPLAY; mkdir -p $(BUILDDIR)/$* && cp -a $< $(BUILDDIR)/$* && cd $(BUILDDIR)/$* && $(AOC) -c $(AOCOFLAGS) $< && cd - && cp -a $(BUILDDIR)/$*/$@ .)
+
+%.aocr: %.aoco
+ (unset DISPLAY; cp -a $< $(BUILDDIR)/$* && cd $(BUILDDIR)/$* && $(AOC) -rtl $(AOCRFLAGS) $< && cd - && cp -a $(BUILDDIR)/$*/$@ .)
+
+%.aocx: %.aocr
+ (unset DISPLAY; cp -a $< $(BUILDDIR)/$* && cd $(BUILDDIR)/$* && $(AOC) $(AOCXFLAGS) $< && cd - && cp -a $(BUILDDIR)/$*/$@ .)
+
+%.sof: %.aocx
+ (unset DISPLAY; cp -a $(BUILDDIR)/$*/flat.sof ./$@)
+
+%.rbf: %.sof
+ (unset DISPLAY; cp -a $(BUILDDIR)/$*/flat.rbf ./$@)
+
+
+%.build:
+ test -f $@ || test -f /tmp/stop || (echo `hostname` && cp `basename $* _$(lastword $(subst _, ,$*))`.cl $*.cl && SEED=$(lastword $(subst _, ,$*)) time make -j1 $*.aocx && fgrep MHz $(BUILDDIR)/$*/$*/quartus_sh_compile.log|tail -n 1) >$@ 2>&1
+
+
diff --git a/applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/inc/common.h b/applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/inc/common.h
new file mode 100644
index 0000000000000000000000000000000000000000..d16ade9618880b8bba4bfcbd42b218f6efb69773
--- /dev/null
+++ b/applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/inc/common.h
@@ -0,0 +1,48 @@
+#include <iostream>
+#include <sstream>
+#include <fstream>
+#include <iomanip>
+
+#define CL_HPP_ENABLE_EXCEPTIONS
+#define CL_HPP_MINIMUM_OPENCL_VERSION 120
+#define CL_HPP_TARGET_OPENCL_VERSION 120
+#define CL_HPP_ENABLE_PROGRAM_CONSTRUCTION_FROM_ARRAY_COMPATIBILITY
+#include <CL/cl2.hpp>
+
+void init(
+ cl::Context &context,
+ std::vector<cl::Device> &devices);
+
+void print_platform(
+ cl::Platform &platform);
+
+void print_device(
+ cl::Device &device,
+ bool marker = false);
+
+std::string get_source(
+ std::string& filename);
+
+std::string get_flags();
+
+cl::Program compile_program(
+ cl::Context& context,
+ cl::Device& device,
+ std::string& source);
+
+void write_source(
+ std::string& source,
+ std::string& filename);
+
+cl::Program get_program(
+ cl::Context& context,
+ cl::Device& device,
+ std::string& filename);
+
+cl::Kernel get_kernel(
+ cl::Program& program,
+ std::string& name);
+
+double compute_runtime(
+ cl::Event& start,
+ cl::Event& end);
diff --git a/applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/readme.css b/applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/readme.css
new file mode 100644
index 0000000000000000000000000000000000000000..ce1c649289c93957c5eeefe2dec8a7b9d8b7d36a
--- /dev/null
+++ b/applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/readme.css
@@ -0,0 +1,261 @@
+/*
+Copyright (C) 2013-2018 Altera Corporation, San Jose, California, USA. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy of this
+software and associated documentation files (the "Software"), to deal in the Software
+without restriction, including without limitation the rights to use, copy, modify, merge,
+publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to
+whom the Software is furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in all copies or
+substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+OTHER DEALINGS IN THE SOFTWARE.
+
+This agreement shall be governed in all respects by the laws of the State of California and
+by the laws of the United States of America.
+*/
+
+body {
+ margin: 0 1em 1em 1em;
+ font-family: sans-serif;
+}
+ul {
+ list-style-type: square;
+}
+pre, code, kbd, samp, tt {
+ font-family: monospace, sans-serif;
+ font-size: 1em;
+}
+
+h1 {
+ font-size: 200%;
+ color: #fff;
+ background-color: #0067a6;
+ margin: 0 -0.5em;
+ padding: 0.25em 0.5em;
+}
+h1 .preheading {
+ font-size: 40%;
+ font-weight: normal;
+}
+h2 {
+ font-size: 125%;
+ background-color: #bae5ff;
+ margin: 1.5em -0.8em 0 -0.8em;
+ padding: 0.2em 0.8em;
+}
+h3 {
+ margin-top: 1.5em;
+ font-size: 100%;
+ border-bottom: 1px dotted #000;
+}
+
+table {
+ border: 2px solid #0067a6;
+ border-collapse: collapse;
+}
+th {
+ border-bottom: 1px solid #0067a6;
+ border-left: 1px dotted #0067a6;
+ border-right: 1px dotted #0067a6;
+ background-color: #bae5ff;
+ padding: 0.3em;
+ font-size: 90%;
+}
+td {
+ padding: 0.3em;
+ border: 1px dotted #0067a6;
+}
+
+table.reqs {
+ margin: 0 auto;
+}
+table.reqs td {
+ white-space: nowrap;
+ text-align: center;
+}
+table.reqs td:first-child,
+table.reqs tr:first-child th:first-child {
+ text-align: left;
+}
+table.reqs td.req {
+ background-color: #b3ef71;
+ font-size: 150%;
+ padding: 0 0.3em;
+}
+table.reqs td.req .either {
+ font-size: 50%;
+}
+table.reqs td.unsupported {
+ white-space: normal;
+ background-color: #ccc;
+ max-width: 20em;
+}
+table.reqs a.note {
+ text-decoration: none;
+}
+ol.req-notes > li {
+ margin-bottom: 0.75em;
+}
+
+table.history {
+ margin: 0 auto;
+}
+table.history td {
+ text-align: center;
+ vertical-align: top;
+}
+table.history .changes {
+ text-align: left;
+}
+table.history tbody tr:first-child td {
+ background-color: #b3ef71;
+}
+table.history ul {
+ margin: 0;
+ padding-left: 1em;
+}
+
+table.pkg-contents {
+ margin: 0 auto;
+}
+table.pkg-contents th,
+table.pkg-contents td {
+ text-align: left;
+ vertical-align: top;
+}
+table.pkg-contents td.path {
+ font-family: monospace, sans-serif;
+ font-size: 1em;
+}
+table.pkg-contents tr.highlight td {
+ background-color: #ffc;
+ font-weight: bold;
+ color: #000;
+}
+table.pkg-contents td p:first-child {
+ margin-top: 0;
+}
+table.pkg-contents td p:last-child {
+ margin-bottom: 0;
+}
+
+table.parameters {
+ margin-left: 3em;
+ margin-right: 3em;
+ font-family: monospace, sans-serif;
+ font-size: 1em;
+}
+table.parameters th,
+table.parameters td {
+ font-family: sans-serif;
+ text-align: center;
+ vertical-align: top;
+}
+table.parameters .name,
+table.parameters .desc {
+ text-align: left;
+}
+table.parameters .name {
+ white-space: nowrap;
+}
+table.parameters td.name,
+table.parameters td.default {
+ font-family: monospace, sans-serif;
+ font-size: 1em;
+}
+table.parameters ul {
+ margin-top: 0;
+}
+table.parameters td ul:last-child {
+ margin-bottom: 0;
+}
+
+table.indent {
+ margin-left: 3em;
+}
+
+.doc .title {
+ background-color: #eee;
+ padding: 0.35em;
+ margin-bottom: 0.5em;
+}
+.doc .title a {
+ font-weight: bold;
+}
+.doc .desc {
+ margin-left: 2em;
+ margin-right: 2em;
+}
+
+.left {
+ text-align: left;
+}
+.center {
+ text-align: center;
+}
+.right {
+ text-align: right;
+}
+
+.mono {
+ font-family: monospace, sans-serif;
+ font-size: 1em;
+}
+.highlight {
+ font-weight: bold;
+ color: #0067a6;
+}
+.nowrap {
+ white-space: nowrap;
+}
+
+.command {
+ font-family: monospace, sans-serif;
+ font-size: 1em;
+ margin: 0 3em;
+ background-color: #ffc;
+ border: 1px solid #aaa;
+ padding: 0.5em 1em;
+}
+.console-output,
+.code-block {
+ display: block;
+ font-family: monospace, sans-serif;
+ font-size: 1em;
+ margin: 0 3em;
+ background-color: #fff;
+ border: 1px solid #aaa !important;
+ padding: 1.8em 1em 0.5em 1em !important;
+ position: relative;
+}
+.console-output .heading,
+.code-block .heading {
+ position: absolute;
+ left: 0;
+ top: 0;
+ width: 100%;
+ font-size: 80%;
+ text-transform: uppercase;
+ background-color: #e8e8e8;
+ padding: 0.3125em 0;
+ border-bottom: 1px dotted #888;
+}
+.console-output .heading span,
+.code-block .heading span {
+ padding: 0 1.25em;
+}
+.not-released {
+ font-weight: bold;
+ color: red;
+}
+.license,
+.trademark {
+ font-size: 80%;
+}
diff --git a/applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/src/common.cpp b/applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/src/common.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..928b8534b95239c6fa0a29f27640984e5605de17
--- /dev/null
+++ b/applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/src/common.cpp
@@ -0,0 +1,189 @@
+#include "common.h"
+
+using namespace std;
+
+ostream &os = clog;
+
+void init(
+ cl::Context &context,
+ vector<cl::Device> &devices)
+{
+ vector<cl::Platform> platforms;
+ cl::Platform::get(&platforms);
+
+ // The selected device
+ int i = 0;
+ const char *platform_name = getenv("PLATFORM");
+
+ if (platform_name == 0)
+ platform_name = getenv("CL_CONTEXT_EMULATOR_DEVICE_INTELFPGA") ? "Intel(R) FPGA Emulation Platform for OpenCL(TM)" : "Intel(R) FPGA SDK for OpenCL(TM)";
+
+ os << ">>> OpenCL environment: " << endl;
+
+ // Iterate all platforms
+ for (cl::Platform &platform : platforms) {
+ print_platform(platform);
+ bool selected = platform.getInfo<CL_PLATFORM_NAME>() == platform_name;
+
+ // Get devices for the current platform
+ vector<cl::Device> devices_;
+ platform.getDevices(CL_DEVICE_TYPE_ALL, &devices_);
+
+ // Iterate all devices
+ for (cl::Device &device : devices_) {
+ if (true)//(selected)
+ devices.push_back(device);
+
+ print_device(device, selected);
+ i++;
+ }
+ }
+ os << endl;
+
+ if (devices.size() == 0) {
+ cerr << "Could not find any device in platform " << platform_name << endl;
+ exit(EXIT_FAILURE);
+ }
+
+ context = cl::Context(devices);
+}
+
+void print_platform(
+ cl::Platform &platform)
+{
+ os << ">>> Platform: " << endl;
+ os << "Name : " << platform.getInfo<CL_PLATFORM_NAME>() << endl;
+ os << "Version : " << platform.getInfo<CL_PLATFORM_VERSION>() << endl;
+ os << "Extensions : " << platform.getInfo<CL_PLATFORM_EXTENSIONS>() << endl;
+ os << endl;
+}
+
+void print_device(
+ cl::Device &device,
+ bool marker)
+{
+ os << ">>> Device: ";
+ if (marker) os << " (selected)";
+ os << endl;
+ os << "Name : " << device.getInfo<CL_DEVICE_NAME>() << endl;
+ os << "Driver version : " << device.getInfo<CL_DRIVER_VERSION>() << endl;
+ os << "Device version : " << device.getInfo<CL_DEVICE_VERSION>() << endl;
+ os << "Compute units : " << device.getInfo<CL_DEVICE_MAX_COMPUTE_UNITS>() << endl;
+ os << "Clock frequency : " << device.getInfo<CL_DEVICE_MAX_CLOCK_FREQUENCY>() << " MHz" << endl;
+ os << "Global memory : " << device.getInfo<CL_DEVICE_GLOBAL_MEM_SIZE>() * 1e-9 << " Gb" << endl;
+ os << "Local memory : " << device.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>() * 1e-6 << " Mb" << endl;
+ os << endl;
+}
+
+string get_source(
+ string& filename)
+{
+ // Source directory
+ string srcdir = "./cl";
+
+ // All helper files to include in build
+ vector<string> helper_files;
+ helper_files.push_back("types.cl");
+ helper_files.push_back("math.cl");
+
+ // Store helper files in string
+ stringstream source_helper_;
+
+ for (int i = 0; i < helper_files.size(); i++) {
+ // Get source filename
+ stringstream source_file_name_;
+ source_file_name_ << srcdir << "/" << helper_files[i];
+ string source_file_name = source_file_name_.str();
+
+ // Read source from file
+ ifstream source_file(source_file_name.c_str());
+ string source(istreambuf_iterator<char>(source_file),
+ (istreambuf_iterator<char>()));
+ source_file.close();
+
+ // Update source helper stream
+ source_helper_ << source;
+ }
+
+ string source_helper = source_helper_.str();
+
+ // Get source filename
+ stringstream source_file_name_;
+ source_file_name_ << srcdir << "/" << filename;
+ string source_file_name = source_file_name_.str();
+
+ // Read kernel source from file
+ ifstream source_file(source_file_name.c_str());
+ string source_kernel(
+ istreambuf_iterator<char>(source_file),
+ (istreambuf_iterator<char>()));
+ source_file.close();
+
+ // Construct full source file
+ stringstream full_source;
+ full_source << source_helper;
+ full_source << source_kernel;
+
+ return full_source.str();
+}
+
+string get_flags()
+{
+ return string("-cl-fast-relaxed-math");
+}
+
+void write_source(
+ string& source,
+ string& filename)
+{
+ cout << ">>> Writing source to: " << filename << endl
+ << endl;
+ ofstream source_output;
+ source_output.open(filename, ofstream::out);
+ source_output << source;
+ source_output.close();
+}
+
+cl::Program get_program(
+ cl::Context& context,
+ cl::Device& device,
+ string& filename)
+{
+ os << ">>> Loading program from binary: " << filename << endl;
+ try {
+ ifstream ifs(filename, ios::in | ios::binary);
+ string str((istreambuf_iterator<char>(ifs)), istreambuf_iterator<char>());
+ cl::Program::Binaries binaries(1, std::make_pair(str.c_str(), str.length()));
+ vector<cl::Device> devices;
+ devices.push_back(device);
+ os << endl;
+ return cl::Program(context, devices, binaries);
+ } catch (cl::Error& error) {
+ cerr << "Loading binary failed: " << error.what() << endl;
+ exit(EXIT_FAILURE);
+ }
+}
+
+cl::Kernel get_kernel(
+ cl::Program& program,
+ string& name)
+{
+ os << ">>> Loading kernel: " << name << endl;
+ try {
+ os << endl;
+ return cl::Kernel(program, name.c_str());
+ } catch (cl::Error& error) {
+ cerr << "Loading kernel failed: " << error.what() << endl;
+ exit(EXIT_FAILURE);
+ }
+}
+
+double compute_runtime(
+ cl::Event& start,
+ cl::Event& end)
+{
+ double runtime = 0;
+ runtime -= start.getProfilingInfo<CL_PROFILING_COMMAND_START>();
+ runtime += end.getProfilingInfo<CL_PROFILING_COMMAND_START>();
+ return runtime * 1e-9;
+}
diff --git a/applications/ta2/designs/ta2_unb2b_mm_demo/host/src/main.cpp b/applications/ta2/designs/ta2_unb2b_mm_demo/host/src/main.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..4eca458dfffc455c1abff5bb81b16f76bf1dd952
--- /dev/null
+++ b/applications/ta2/designs/ta2_unb2b_mm_demo/host/src/main.cpp
@@ -0,0 +1,111 @@
+/* *************************************************************************
+* Copyright 2020
+* ASTRON (Netherlands Institute for Radio Astronomy) <http://www.astron.nl/>
+* P.O.Box 2, 7990 AA Dwingeloo, The Netherlands
+*
+* Licensed under the Apache License, Version 2.0 (the "License");
+* you may not use this file except in compliance with the License.
+* You may obtain a copy of the License at
+*
+* http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing, software
+* distributed under the License is distributed on an "AS IS" BASIS,
+* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+* See the License for the specific language governing permissions and
+* limitations under the License.
+* *********************************************************************** */
+
+/* *************************************************************************
+* Author:
+* . Reinier vd Walle
+* Purpose:
+* . Test the ta2_unb2b_mm_demo OpenCL application in emulator
+* Description:
+* . Run: -> make ta2_unb2b_mm_demo
+* . Navigate to -> cd $RADIOHDL_WORK/unb2b/OpenCL/ta2_unb2b_mm_demo/bin
+* . Execute -> CL_CONTEXT_EMULATOR_DEVICE_INTELFPGA=1 ./host
+* *********************************************************************** */
+#include <CL/cl_ext_intelfpga.h>
+#include <iostream>
+#include <fstream>
+#include <vector>
+#include "common.h"
+#include <unistd.h>
+
+using namespace std;
+int main(int argc, char **argv)
+{
+ if (argc > 2) {
+ cerr << "usage: " << argv[0] << " [ta2_unb2b_mm_demo.aocx]" << endl;
+ exit(1);
+ }
+
+ // Initialize OpenCL
+ cl::Context context;
+ vector<cl::Device> devices;
+ init(context, devices);
+ cl::Device &device = devices[0];
+
+ // Get program
+ string filename_bin = string(argc == 2 ? argv[1] : "ta2_unb2b_mm_demo.aocx");
+ cl::Program program = get_program(context, device, filename_bin);
+
+
+ // Setup command queues
+ vector<cl::CommandQueue> queues(4);
+
+ for (cl::CommandQueue &queue : queues) {
+ queue = cl::CommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE);
+ }
+
+ cl::Event computeDone[4];
+
+ // Setup FPGA kernels
+ cl::Kernel mmInController(program, "mm_in_controller");
+ cl::Kernel mmOutController(program, "mm_out_controller");
+ cl::Kernel processA(program, "process_a");
+ cl::Kernel processB(program, "process_b");
+
+
+ // Run FPGA kernels
+ clog << ">>> Run fpga" << endl;
+ try {
+ queues[0].enqueueTask(processA, nullptr, &computeDone[0]);
+ queues[1].enqueueTask(processB, nullptr, &computeDone[1]);
+ queues[2].enqueueTask(mmOutController, nullptr, &computeDone[2]);
+ queues[3].enqueueTask(mmInController, nullptr, &computeDone[3]);
+
+ } catch (cl::Error &error) {
+ cerr << "Error launching kernel: " << error.what() << endl;
+ exit(EXIT_FAILURE);
+ }
+
+ // Write IO channel file
+ vector<char> cmdVecs[] = {{'A', 'B', 'C', 'D', 0x33, 0x00, 0x00, 0x00, 0x01}, //write on undefined address
+ {'E', 'F', 'G', 'H', 0x00, 0x00, 0x00, 0x00, 0x01}, // write on addr 0
+ {'I', 'J', 'K', 'L', 0x01, 0x00, 0x00, 0x00, 0x01}, // write on addr 1
+ {'M', 'N', 'O', 'P', 0x02, 0x00, 0x00, 0x00, 0x01}, // write on addr 2
+ {0x00, 0x00, 0x00, 0x00, 0x33, 0x00, 0x00, 0x00, 0x00}, // read on undefined address
+ {0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00}, // read on addr 0
+ {0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00}, // read on addr 1
+ {0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x00}}; // read on addr 2
+
+ ofstream output_fileA("kernel_input_mm");
+ ostream_iterator<char> output_iteratorA(output_fileA, "");
+ for (int i = 0; i < 8; i++)
+ copy(cmdVecs[i].begin(), cmdVecs[i].end(), output_iteratorA);
+
+ output_fileA.close();
+ clog << ">>> Written IO file" << endl;
+
+ // wait for mm_out_controller to be finished
+ computeDone[2].wait();
+
+ // print output IO channel file
+ const string inputFileB = "kernel_output_mm";
+ ifstream fileB(inputFileB);
+ clog << fileB.rdbuf() << endl;
+
+ return EXIT_SUCCESS;
+}
diff --git a/applications/ta2/designs/ta2_unb2b_mm_demo/ta2_unb2b_mm_demo.cl b/applications/ta2/designs/ta2_unb2b_mm_demo/ta2_unb2b_mm_demo.cl
new file mode 100644
index 0000000000000000000000000000000000000000..77b2c03ff8f5947c91ae061d1796101ea6f38493
--- /dev/null
+++ b/applications/ta2/designs/ta2_unb2b_mm_demo/ta2_unb2b_mm_demo.cl
@@ -0,0 +1,199 @@
+/* *************************************************************************
+* Copyright 2020
+* ASTRON (Netherlands Institute for Radio Astronomy) <http://www.astron.nl/>
+* P.O.Box 2, 7990 AA Dwingeloo, The Netherlands
+*
+* Licensed under the Apache License, Version 2.0 (the "License");
+* you may not use this file except in compliance with the License.
+* You may obtain a copy of the License at
+*
+* http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing, software
+* distributed under the License is distributed on an "AS IS" BASIS,
+* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+* See the License for the specific language governing permissions and
+* limitations under the License.
+* *********************************************************************** */
+
+/* *************************************************************************
+* Author:
+* . Reinier vd Walle
+* Purpose:
+* . Demonstrate Monitor and Control interface
+* Description:
+* . This application implements a way to use the MM IO channels.
+* *********************************************************************** */
+
+#pragma OPENCL EXTENSION cl_intel_channels : enable
+
+#include <ihc_apint.h>
+
+struct mm_in {
+ uint wrdata;
+ uint address;
+ uchar wr;
+} __attribute__((packed));
+
+struct mm_out {
+ uint rddata;
+} __attribute__((packed));
+
+
+channel struct mm_in ch_in_mm __attribute__((depth(0))) __attribute__((io("kernel_input_mm")));
+channel struct mm_out ch_out_mm __attribute__((depth(0))) __attribute__((io("kernel_output_mm")));
+
+
+struct reg {
+ uint offset;
+ uint size;
+} __attribute__((packed));
+
+#define REGISTER_A 0x00
+#define REGISTER_B 0x02
+
+#define NR_MM 2
+channel struct mm_in mm_channel_in[NR_MM] __attribute__((depth(0)));
+channel struct mm_out mm_channel_out[NR_MM+1] __attribute__((depth(0))); // 1 extra channel for undefined addresses
+
+__attribute__((max_global_work_dim(0)))
+#ifndef EMULATOR
+__attribute__((autorun))
+#endif
+__kernel void mm_in_controller()
+{
+ const struct reg regmap[NR_MM] = {
+ {REGISTER_A, 2},
+ {REGISTER_B, 2}
+ };
+ while(1)
+ {
+ bool undefined = true;
+ struct mm_in mm_request = read_channel_intel(ch_in_mm);
+ #pragma unroll
+ for (int i = 0; i < NR_MM; i++)
+ {
+ if (mm_request.address >= regmap[i].offset && mm_request.address < (regmap[i].offset + regmap[i].size))
+ {
+ undefined = false;
+ struct mm_in local_mm_request;
+ local_mm_request.wr = mm_request.wr;
+ local_mm_request.wrdata = mm_request.wrdata;
+ local_mm_request.address = mm_request.address - regmap[i].offset;
+ write_channel_intel(mm_channel_in[i], local_mm_request);
+ }
+ }
+
+ if (undefined && mm_request.wr == 0) { // undefined address
+ struct mm_out zero_response;
+ zero_response.rddata = 0;
+ write_channel_intel(mm_channel_out[NR_MM], zero_response);
+ }
+
+ }
+}
+
+__attribute__((max_global_work_dim(0)))
+#ifndef EMULATOR
+__attribute__((autorun))
+#endif
+__kernel void mm_out_controller()
+{
+#ifdef EMULATOR
+ for(int x = 0; x < 4; )
+#else
+ while(1)
+#endif
+ {
+ struct mm_out mm_response;
+ for (int i = 0; i < NR_MM+1; i++)
+ {
+ bool valid;
+ mm_response = read_channel_nb_intel(mm_channel_out[i], &valid);
+ if (valid)
+ {
+ write_channel_intel(ch_out_mm, mm_response);
+#ifdef EMULATOR
+ x++;
+#endif
+ }
+ }
+ }
+}
+
+
+__attribute__((max_global_work_dim(0)))
+#ifndef EMULATOR
+__attribute__((autorun))
+#endif
+__kernel void process_a()
+{
+ uint keep = 0; //address 0, value is stored when written
+ uint acc = 0; //address 1, value is written value + 1
+ while(1){
+ struct mm_in mm_request = read_channel_intel(mm_channel_in[0]);
+ struct mm_out mm_response;
+ if (0 == mm_request.address){
+ if(mm_request.wr > 0) //write request
+ {
+ keep = mm_request.wrdata;
+ } else { //read request
+ mm_response.rddata = keep;
+ }
+ }
+
+ if (1 == mm_request.address){
+ if(mm_request.wr > 0) //write request
+ {
+ acc = mm_request.wrdata+1;
+ } else { //read request
+ mm_response.rddata = acc;
+ }
+ }
+
+ if(mm_request.wr == 0)
+ write_channel_intel(mm_channel_out[0], mm_response);
+ }
+}
+
+
+__attribute__((max_global_work_dim(0)))
+#ifndef EMULATOR
+__attribute__((autorun))
+#endif
+__kernel void process_b()
+{
+ uint keep = 0; //address 0, value is stored when written
+ uint acc = 0; //address 1, value is written value + 2
+ while(1){
+ struct mm_in mm_request = read_channel_intel(mm_channel_in[1]);
+ struct mm_out mm_response;
+ if (0 == mm_request.address){
+ if(mm_request.wr > 0) //write request
+ {
+ keep = mm_request.wrdata;
+ } else { //read request
+ mm_response.rddata = keep;
+ }
+ }
+
+ if (1 == mm_request.address){
+ if(mm_request.wr > 0) //write request
+ {
+ acc = mm_request.wrdata+2;
+ } else { //read request
+ mm_response.rddata = acc;
+ }
+ }
+
+ if(mm_request.wr == 0)
+ write_channel_intel(mm_channel_out[1], mm_response);
+ }
+}
+
+
+__attribute__((max_global_work_dim(0)))
+__kernel void dummy()
+{
+}
+