Initial commit of ta2_unb2b_mm_demo OpenCL application.

applications/ta2/designs/ta2_unb2b_mm_demo/Makefile

+######################
+###     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
+
+

applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/inc/common.h

+#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);

applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/readme.css

+/*
+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%;
+}

applications/ta2/designs/ta2_unb2b_mm_demo/host/lib/common/src/common.cpp

+#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;
+}

applications/ta2/designs/ta2_unb2b_mm_demo/host/src/main.cpp

+/* *************************************************************************
+* 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;
+}

applications/ta2/designs/ta2_unb2b_mm_demo/ta2_unb2b_mm_demo.cl

+/* *************************************************************************
+* 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()
+{
+}
+