From 6ee6177c9ec87e0d2a145c79b96a16977ddcc0c6 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Marcel=20L=C3=BCtke=20Dreimann?=
 <marcel.luetkedreimann@uos.de>
Date: Thu, 30 Jun 2022 12:40:56 +0200
Subject: [PATCH] added gpgpu driver

---
 .gitmodules                                   |   3 +
 repos/dde_uos-intel-gpgpu/README              |   0
 repos/dde_uos-intel-gpgpu/include/gpgpu.h     |   5 +
 repos/dde_uos-intel-gpgpu/run/gpgpu.run       |  67 ++++++
 .../src/gpgpu/gpgpu_genode.cc                 | 136 ++++++++++++
 .../src/gpgpu/gpgpu_genode.h                  | 135 ++++++++++++
 repos/dde_uos-intel-gpgpu/src/gpgpu/main.cc   |  82 +++++++
 repos/dde_uos-intel-gpgpu/src/gpgpu/stubs.cc  |  67 ++++++
 repos/dde_uos-intel-gpgpu/src/gpgpu/target.mk |  20 ++
 repos/dde_uos-intel-gpgpu/src/gpgpu/test.cc   | 206 ++++++++++++++++++
 repos/dde_uos-intel-gpgpu/src/gpgpu/test.h    |  10 +
 repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu |   1 +
 12 files changed, 732 insertions(+)
 create mode 100644 .gitmodules
 create mode 100644 repos/dde_uos-intel-gpgpu/README
 create mode 100644 repos/dde_uos-intel-gpgpu/include/gpgpu.h
 create mode 100644 repos/dde_uos-intel-gpgpu/run/gpgpu.run
 create mode 100644 repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.cc
 create mode 100644 repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.h
 create mode 100644 repos/dde_uos-intel-gpgpu/src/gpgpu/main.cc
 create mode 100644 repos/dde_uos-intel-gpgpu/src/gpgpu/stubs.cc
 create mode 100644 repos/dde_uos-intel-gpgpu/src/gpgpu/target.mk
 create mode 100644 repos/dde_uos-intel-gpgpu/src/gpgpu/test.cc
 create mode 100644 repos/dde_uos-intel-gpgpu/src/gpgpu/test.h
 create mode 160000 repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu

diff --git a/.gitmodules b/.gitmodules
new file mode 100644
index 0000000000..84f727fe1d
--- /dev/null
+++ b/.gitmodules
@@ -0,0 +1,3 @@
+[submodule "repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu"]
+	path = repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu
+	url = git@ess-git.inf.uos.de:software/uos-intel-gpgpu.git
diff --git a/repos/dde_uos-intel-gpgpu/README b/repos/dde_uos-intel-gpgpu/README
new file mode 100644
index 0000000000..e69de29bb2
diff --git a/repos/dde_uos-intel-gpgpu/include/gpgpu.h b/repos/dde_uos-intel-gpgpu/include/gpgpu.h
new file mode 100644
index 0000000000..2f0e46d8b9
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/include/gpgpu.h
@@ -0,0 +1,5 @@
+// driver include
+#include "../src/uos-intel-gpgpu/driver/gpgpu_driver.h"
+
+// include for genode wrapper
+#include "../src/gpgpu/gpgpu_genode.h"
diff --git a/repos/dde_uos-intel-gpgpu/run/gpgpu.run b/repos/dde_uos-intel-gpgpu/run/gpgpu.run
new file mode 100644
index 0000000000..60484570e5
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/run/gpgpu.run
@@ -0,0 +1,67 @@
+# build config
+set build_components { core init gpgpu timer }
+
+# platform config
+set use_acpica_as_acpi_drv 0
+source ${genode_dir}/repos/base/run/platform_drv.inc
+proc platform_drv_policy {} {
+	global use_acpica_as_acpi_drv
+	set policy ""
+
+	append_if $use_acpica_as_acpi_drv policy {
+			<policy label="acpi_drv -> ">    <pci class="ALL"/> </policy>}
+
+	append policy {
+			<policy label_prefix="gpgpu"> <pci class="ALL"/> </policy>}
+
+	return $policy
+}
+append_platform_drv_build_components
+build $build_components
+
+# boot dir
+create_boot_directory
+
+# other config
+append config {
+    <config>
+        <parent-provides>
+            <service name="ROM"/>
+            <service name="IRQ"/>
+            <service name="IO_MEM"/>
+            <service name="IO_PORT"/>
+            <service name="PD"/>
+            <service name="RM"/>
+            <service name="CPU"/>
+            <service name="LOG"/>
+        </parent-provides>
+        <default-route>
+            <any-service> <parent/> <any-child/> </any-service>
+        </default-route>
+        <default caps="100"/>
+        <start name="timer">
+            <resource name="RAM" quantum="10M"/>
+            <provides><service name="Timer"/></provides>
+    	</start>
+}
+
+append_platform_drv_config
+
+append config {
+    <start name="gpgpu">
+        <resource name="RAM" quantum="32M"/>
+    </start>
+</config>}
+
+install_config $config
+
+# boot modules
+set boot_modules {
+	core ld.lib.so init gpgpu timer
+}
+append_platform_drv_boot_modules
+build_boot_image $boot_modules
+
+# qemu stuff
+append qemu_args "-nographic -m 256"
+run_genode_until {This is the UOS Intel GPGPU End!.*\n} 10
diff --git a/repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.cc b/repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.cc
new file mode 100644
index 0000000000..0bb6ba702c
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.cc
@@ -0,0 +1,136 @@
+#include "gpgpu_genode.h"
+
+#define GENODE // use genodes stdint header
+#include "../uos-intel-gpgpu/driver/gpgpu_driver.h"
+
+void gpgpu_genode::handleInterrupt()
+{
+    // handle the gpu interrupt
+    GPGPU_Driver& gpgpudriver = GPGPU_Driver::getInstance();
+    gpgpudriver.handleInterrupt();
+    gpgpudriver.runNext();
+
+    // ack the irq
+    irq->ack_irq();
+}
+
+gpgpu_genode::gpgpu_genode(Env& e) : env(e), heap{ e.ram(), e.rm() }, alloc(&heap), ram_cap(), mapped_base(0), base(0), pci(e), dev(), prev_dev(), irq(nullptr), dispatcher(env.ep(), *this, &gpgpu_genode::handleInterrupt)
+{
+    // size of avaible memory for allocator
+    const unsigned long size = 0x1000 * 0x1000;
+
+    // allocate chunk of ram
+    ram_cap = e.ram().alloc(size);
+    mapped_base = e.rm().attach(ram_cap);
+    base = Dataspace_client(ram_cap).phys_addr();
+
+    // use this ram for allocator
+    alloc.add_range(mapped_base, size);
+}
+
+gpgpu_genode::~gpgpu_genode()
+{
+    // release pci dev and free allocator memory
+    pci.release_device(dev);
+    env.ram().free(ram_cap);
+}
+
+void* gpgpu_genode::aligned_alloc(uint32_t alignment, uint32_t size)
+{
+    void* ptr;
+    alloc.alloc_aligned(size, &ptr, alignment);
+    return ptr;
+}
+
+void gpgpu_genode::free(void* addr)
+{
+    alloc.free(addr);
+}
+
+void gpgpu_genode::createPCIConnection(uint8_t bus, uint8_t device, uint8_t function)
+{
+    // get first device
+    pci.with_upgrade([&] () { dev = pci.first_device(); });
+
+    while (dev.valid()) {
+        // release old one
+        pci.release_device(prev_dev);
+        prev_dev = dev;
+
+        // get next one
+        pci.with_upgrade([&] () { dev = pci.next_device(dev); });
+
+        // check if this is the right one
+        Platform::Device_client client(dev);
+        uint8_t b, d, f;
+        client.bus_address(&b, &d, &f);
+        if(b == bus && d == device && f == function)
+        {
+            break;
+        }
+    }
+
+    // we did not find the right one
+    if (!dev.valid())
+    {
+        Genode::error("[GENODE_GPGPU]: Could not find PCI dev: ", bus, device, function);
+        return;
+    }
+}
+
+uint32_t gpgpu_genode::readPCI(uint32_t addr)
+{
+    Platform::Device_client client(dev);
+    return client.config_read(addr, Platform::Device::ACCESS_32BIT);
+}
+
+void gpgpu_genode::writePCI(uint32_t addr, uint32_t val)
+{
+    Platform::Device_client client(dev);
+    pci.with_upgrade([&] () {
+        client.config_write(addr, val, Platform::Device::ACCESS_32BIT);
+    });
+}
+
+addr_t gpgpu_genode::getVirtBarAddr(uint8_t bar_id) const
+{
+    // get virt bar id (why does this exist?)
+    Platform::Device_client dc(dev);
+    Platform::Device::Resource res = dc.resource(bar_id);
+    uint8_t genodeBarID = dc.phys_bar_to_virt(bar_id);
+
+    // create io mem session
+    Genode::Io_mem_session_capability cap = dc.io_mem(genodeBarID);
+    if (!cap.valid())
+    {
+        Genode::error("[GENODE_GPGPU]: IO memory session is not valid");
+        return 0;
+    }
+
+    // get dataspace cap
+    Genode::Io_mem_session_client mem(cap);
+    Genode::Io_mem_dataspace_capability mem_ds(mem.dataspace());
+    if (!mem_ds.valid())
+    {
+        Genode::error("[GENODE_GPGPU]: IO mem dataspace cap not valid");
+        return 0;
+    }
+
+    // add addr to rm and get virt addr
+    addr_t vaddr = env.rm().attach(mem_ds);
+    vaddr |= res.base() & 0xfff;
+    return vaddr;
+}
+
+void gpgpu_genode::registerInterruptHandler()
+{
+    Platform::Device_client client(dev);
+    static Irq_session_client irq_client(client.irq(0)); // 0 ??
+    irq = &irq_client;
+
+    // set dispatcher
+    irq->sigh(dispatcher);
+
+    // initial ack
+    irq->ack_irq();
+}
diff --git a/repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.h b/repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.h
new file mode 100644
index 0000000000..b448f3e988
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.h
@@ -0,0 +1,135 @@
+#ifndef GPGPU_GENODE_H
+#define GPGPU_GENODE_H
+
+// stdint
+#include <base/fixed_stdint.h>
+using namespace Genode;
+
+// allocator
+#include <base/heap.h>
+#include <base/allocator_avl.h>
+#include <dataspace/client.h>
+
+// pci
+#include <platform_session/connection.h>
+#include <platform_device/client.h>
+#include <io_mem_session/connection.h>
+#include <io_port_session/connection.h>
+
+// interrupts
+#include <irq_session/connection.h>
+
+class gpgpu_genode
+{
+private:
+    // genode enviroment
+    Env& env;
+
+    // allocator
+    Heap heap;
+    Allocator_avl alloc;
+    Ram_dataspace_capability ram_cap;
+    addr_t mapped_base;
+    addr_t base;
+
+    // pci
+	Platform::Connection pci;
+	Platform::Device_capability dev;
+	Platform::Device_capability prev_dev;
+    
+    // interrupts
+    Irq_session_client* irq;
+    Signal_handler<gpgpu_genode> dispatcher;
+
+    // do not allow copies
+    gpgpu_genode(const gpgpu_genode& copy) = delete;
+    gpgpu_genode& operator=(const gpgpu_genode& src) = delete;
+
+    /**
+     * @brief Interrupt handler
+     * 
+     */
+    void handleInterrupt();
+
+public:
+    /**
+     * @brief Construct a new gpgpu genode object
+     * 
+     * @param e 
+     */
+    gpgpu_genode(Env& e);
+
+    /**
+     * @brief Destroy the gpgpu genode object
+     * 
+     */
+    ~gpgpu_genode();
+
+    /**
+     * @brief allocate aligned memory
+     * 
+     * @param alignment the alignment
+     * @param size the size in bytes
+     * @return void* the address of the allocated memory
+     */
+    void* aligned_alloc(uint32_t alignment, uint32_t size);
+
+    /**
+     * @brief free memory
+     * 
+     * @param addr the address of the memory to be freed
+     */
+    void free(void* addr);
+
+    /**
+     * @brief converts a virtual address into a physical address
+     * 
+     * @param virt the virtual address
+     * @return addr_t the physical address
+     */
+    addr_t virt_to_phys(addr_t virt) const
+    {
+		return virt - mapped_base + base;
+    }
+
+    /**
+     * @brief creates a connection to the PCI device. This has to be called before any read/write to the PCI device!
+     * 
+     * @param bus the bus id
+     * @param device the device id
+     * @param function the function id
+     */
+    void createPCIConnection(uint8_t bus, uint8_t device, uint8_t function);
+
+    /**
+     * @brief read from pci config space
+     * 
+     * @param addr the address to read from
+     * @return uint32_t the value
+     */
+    uint32_t readPCI(uint32_t addr);
+
+    /**
+     * @brief write to pci config space (some register are protected by genode!)
+     * 
+     * @param addr the address to write to
+     * @param val the value to write
+     */
+    void writePCI(uint32_t addr, uint32_t val);
+
+    /**
+     * @brief Get the Virt Bar Addr object
+     * 
+     * @param bar_id 
+     * @return addr_t 
+     */
+    addr_t getVirtBarAddr(uint8_t bar_id) const;
+
+    /**
+     * @brief register the interrupt handler for the current PCI device
+     * 
+     */
+    void registerInterruptHandler();
+};
+
+#endif // GPGPU_GENODE_H
diff --git a/repos/dde_uos-intel-gpgpu/src/gpgpu/main.cc b/repos/dde_uos-intel-gpgpu/src/gpgpu/main.cc
new file mode 100644
index 0000000000..61eb3d78bd
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/src/gpgpu/main.cc
@@ -0,0 +1,82 @@
+#include <base/component.h>
+
+#define GENODE // use genodes stdint header
+#include "../uos-intel-gpgpu/driver/gpgpu_driver.h"
+#include "gpgpu_genode.h"
+
+//#define TEST // test stubs only (works with qemu)
+#ifdef TEST
+#include "../uos-intel-gpgpu/stubs.h"
+#else
+#include "test.h"
+#endif // TEST
+
+gpgpu_genode* _global_gpgpu_genode;
+
+void Component::construct(Genode::Env& e)
+{
+    Genode::log("Hello world: UOS Intel GPGPU!");
+    Genode::log("Build: ", __TIMESTAMP__);
+
+    // init globals
+    static gpgpu_genode gg(e);
+    _global_gpgpu_genode = &gg;
+
+#ifdef TEST
+    // test prink
+    printk("Hello printk: %d", 42);
+
+    // test alloc
+    uint8_t* test = (uint8_t*)uos_aligned_alloc(0x1000, 0x1000);
+    uint64_t addr = (uint64_t)test;
+    if((addr & 0xFFF) != 0)
+    {
+        Genode::error("mem alignment failed: ", addr);
+    }
+    for(int i = 0; i < 0x1000; i++)
+    {
+        test[i] = 0x42;
+    }
+    for(int i = 0; i < 0x1000; i++)
+    {
+        if(test[i] != 0x42)
+        {
+            Genode::error("mem write or read failed!");
+            break;
+        }
+    }
+    free(test);
+    Genode::log("Allocator test finished!");
+
+    // test pci
+    uint32_t base = calculatePCIConfigHeaderAddress(0, 2 , 0);
+    uint32_t dev_ven = readPCIConfigSpace(base + 0);
+    if((dev_ven & 0xFFFF) == 0x8086)
+    {
+        Genode::log("PCI test successful!");
+    }
+    else
+    {
+        Genode::error("PCI test failed!");
+    }
+
+    // test pci memory
+    uint8_t* test2 = (uint8_t*)_global_gpgpu_genode->getVirtBarAddr(0);
+    test2[0x42] = 0x42;
+    Genode::log("PCI memory test finished!");
+
+    // test interrupts
+    _global_gpgpu_genode->registerInterruptHandler();
+    Genode::log("Interrupt test finished!");
+#else
+    // init driver
+    GPGPU_Driver& gpgpudriver = GPGPU_Driver::getInstance();
+    gpgpudriver.init(0);
+    _global_gpgpu_genode->registerInterruptHandler();
+
+    // run the test and hope the best
+    run_gpgpu_test();
+#endif // TEST
+
+    Genode::log("This is the UOS Intel GPGPU End!");
+}
diff --git a/repos/dde_uos-intel-gpgpu/src/gpgpu/stubs.cc b/repos/dde_uos-intel-gpgpu/src/gpgpu/stubs.cc
new file mode 100644
index 0000000000..97036aa517
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/src/gpgpu/stubs.cc
@@ -0,0 +1,67 @@
+// stdint
+#include <base/fixed_stdint.h>
+using namespace Genode;
+
+// printk
+#include <base/log.h>
+#include <base/snprintf.h>
+#include <util/string.h>
+
+// genode instance
+#include "gpgpu_genode.h"
+extern gpgpu_genode* _global_gpgpu_genode;
+
+// printing (optional)
+extern "C" int printk(const char* str, ...)
+{
+    va_list list;
+    va_start(list, str);
+
+    char buff[256];
+	String_console sc(buff, sizeof(buff));
+	sc.vprintf(str, list);
+
+    va_end(list);
+
+    Genode::log("[GPU] ", Genode::Cstring(buff));
+    return 0;
+}
+
+// allocator
+extern "C" void* uos_aligned_alloc(uint32_t alignment, uint32_t size)
+{
+    return _global_gpgpu_genode->aligned_alloc(alignment, size);
+}
+
+extern "C" void free(void* addr)
+{
+    _global_gpgpu_genode->free(addr);
+}
+
+// pci
+extern "C" uint32_t calculatePCIConfigHeaderAddress(uint8_t bus, uint8_t device, uint8_t function)
+{
+    _global_gpgpu_genode->createPCIConnection(bus, device, function);
+	return 0;
+}
+
+extern "C" uint32_t readPCIConfigSpace(uint32_t addr)
+{
+    return _global_gpgpu_genode->readPCI(addr);
+}
+
+extern "C" void writePCIConfigSpace(uint32_t address, uint32_t value)
+{
+    _global_gpgpu_genode->writePCI(address, value);
+}
+
+// address model
+extern "C" void* getVirtBarAddr(uint8_t bar_id)
+{
+    return (void*)_global_gpgpu_genode->getVirtBarAddr(bar_id);
+}
+
+extern "C" void* virt_to_phys(void* addr)
+{
+    return (void*)_global_gpgpu_genode->virt_to_phys((addr_t)addr);
+}
diff --git a/repos/dde_uos-intel-gpgpu/src/gpgpu/target.mk b/repos/dde_uos-intel-gpgpu/src/gpgpu/target.mk
new file mode 100644
index 0000000000..dcb68e06d7
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/src/gpgpu/target.mk
@@ -0,0 +1,20 @@
+TARGET = gpgpu
+REQUIRES = x86_64
+
+SRC_CC = main.cc gpgpu_genode.cc stubs.cc test.cc
+LIBS   = base
+
+UOS_INTEL_GPGPU = uos-intel-gpgpu-link-cxx.o
+EXT_OBJECTS = $(BUILD_BASE_DIR)/bin/$(UOS_INTEL_GPGPU)
+
+$(TARGET): $(UOS_INTEL_GPGPU)
+
+$(UOS_INTEL_GPGPU): $(SRC_CC)
+	$(MSG_BUILD) "Building uos-intel-gpgpu..."
+	$(MAKE) -C $(REP_DIR)/src/uos-intel-gpgpu/
+	cp $(REP_DIR)/src/uos-intel-gpgpu/build/$(UOS_INTEL_GPGPU) $(BUILD_BASE_DIR)/bin/.
+
+clean_uos-intel-gpgpu:
+	$(MAKE) -C $(REP_DIR)/src/uos-intel-gpgpu/ clean
+
+clean: clean_uos-intel-gpgpu
diff --git a/repos/dde_uos-intel-gpgpu/src/gpgpu/test.cc b/repos/dde_uos-intel-gpgpu/src/gpgpu/test.cc
new file mode 100644
index 0000000000..ba7d7f921c
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/src/gpgpu/test.cc
@@ -0,0 +1,206 @@
+#define GENODE // use genodes stdint header
+#include "../uos-intel-gpgpu/driver/gpgpu_driver.h"
+#include "../uos-intel-gpgpu/stubs.h"
+
+#define ELEMENTS 4096
+
+uint32_t* in;
+uint32_t* out;
+
+/*
+kernel void clmain(global const unsigned int* in, global unsigned int* out)
+{
+    unsigned int i = get_global_id(0);
+    out[i] = in[i];
+}
+*/
+
+static unsigned char test_Gen9core_gen[] = {
+  0x43, 0x54, 0x4e, 0x49, 0x2e, 0x04, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
+  0x08, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x4c, 0x04, 0x96, 0x2a, 0x25, 0xad, 0x06, 0x1f,
+  0x99, 0x00, 0x72, 0x8d, 0x08, 0x00, 0x00, 0x00, 0xac, 0x03, 0x00, 0x00,
+  0x80, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00,
+  0x88, 0x00, 0x00, 0x00, 0xe0, 0x00, 0x00, 0x00, 0x63, 0x6c, 0x6d, 0x61,
+  0x69, 0x6e, 0x00, 0x00, 0x01, 0x00, 0x60, 0x00, 0x0c, 0x02, 0x60, 0x20,
+  0x00, 0x00, 0x20, 0x00, 0x00, 0x00, 0x00, 0x00, 0x06, 0x80, 0x00, 0x00,
+  0x04, 0x00, 0x00, 0x30, 0x00, 0x10, 0x00, 0x16, 0xc0, 0x04, 0xc0, 0x04,
+  0x41, 0x00, 0x00, 0x00, 0x2c, 0x0a, 0x80, 0x20, 0x10, 0x01, 0x00, 0x0a,
+  0x64, 0x00, 0x00, 0x00, 0x01, 0x4d, 0x00, 0x20, 0x07, 0x7f, 0x03, 0x00,
+  0x40, 0x00, 0x80, 0x00, 0x28, 0x0a, 0xa0, 0x20, 0x80, 0x00, 0x00, 0x12,
+  0x20, 0x00, 0xb1, 0x00, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x20, 0x21,
+  0x80, 0x00, 0x00, 0x12, 0x40, 0x00, 0xb1, 0x00, 0x40, 0x96, 0x01, 0x20,
+  0x07, 0x05, 0x05, 0x07, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x20, 0x21,
+  0x20, 0x01, 0x8d, 0x0a, 0xe0, 0x00, 0x00, 0x00, 0x09, 0x00, 0x80, 0x00,
+  0x28, 0x0a, 0xa0, 0x20, 0xa0, 0x00, 0x8d, 0x1e, 0x02, 0x00, 0x02, 0x00,
+  0x09, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x20, 0x21, 0x20, 0x01, 0x8d, 0x1e,
+  0x02, 0x00, 0x02, 0x00, 0x31, 0x00, 0x80, 0x0c, 0x68, 0x02, 0x60, 0x21,
+  0xa0, 0x00, 0x00, 0x06, 0x00, 0x5e, 0x20, 0x04, 0x31, 0x20, 0x80, 0x0c,
+  0x68, 0x02, 0xa0, 0x21, 0x20, 0x01, 0x00, 0x06, 0x00, 0x5e, 0x20, 0x04,
+  0x33, 0x00, 0x80, 0x0c, 0x70, 0xb0, 0x00, 0x00, 0xa2, 0x00, 0x00, 0x00,
+  0x01, 0x5e, 0x02, 0x04, 0x33, 0x20, 0x80, 0x0c, 0x70, 0xd0, 0x00, 0x00,
+  0x22, 0x01, 0x00, 0x00, 0x01, 0x5e, 0x02, 0x04, 0x31, 0x00, 0x60, 0x07,
+  0x04, 0x02, 0x00, 0x20, 0xe0, 0x0f, 0x00, 0x06, 0x10, 0x00, 0x00, 0x82,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x82, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
+  0x00, 0xc0, 0xff, 0x83, 0x00, 0x00, 0x00, 0x03, 0x7f, 0x00, 0xff, 0x1f,
+  0x00, 0x00, 0xe0, 0x0f, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0xff, 0x83, 0x00, 0x00, 0x00, 0x03,
+  0x7f, 0x00, 0xff, 0x1f, 0x00, 0x00, 0xe0, 0x0f, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x40, 0x00, 0x00, 0x00, 0x13, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x15, 0x00, 0x00, 0x00, 0x18, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x80, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x14, 0x00, 0x00, 0x00,
+  0x80, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00,
+  0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
+  0x04, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
+  0x28, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x08, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00,
+  0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00,
+  0x04, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
+  0x28, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x14, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00,
+  0x2b, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00,
+  0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
+  0x28, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x30, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x34, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
+  0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00,
+  0x1c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x38, 0x00, 0x00, 0x00,
+  0x04, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x00, 0x00,
+  0x24, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x20, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff,
+  0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x00, 0x00,
+  0x24, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00,
+  0x28, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff,
+  0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00,
+  0x0c, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00,
+  0x44, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
+  0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x17, 0x00, 0x00, 0x00, 0x84, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x20, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x80, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00,
+  0x00, 0x00, 0x00, 0x00, 0x1b, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00,
+  0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
+  0x48, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
+  0x08, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
+  0x08, 0x00, 0x00, 0x00, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c,
+  0x00, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
+  0x69, 0x6e, 0x00, 0x00, 0x75, 0x69, 0x6e, 0x74, 0x2a, 0x3b, 0x38, 0x00,
+  0x63, 0x6f, 0x6e, 0x73, 0x74, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
+  0x48, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
+  0x08, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
+  0x08, 0x00, 0x00, 0x00, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c,
+  0x00, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
+  0x6f, 0x75, 0x74, 0x00, 0x75, 0x69, 0x6e, 0x74, 0x2a, 0x3b, 0x38, 0x00,
+  0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00
+};
+
+void cleanUp()
+{
+    printk("Yeah.. it finished!");
+
+    // set gpu frequency to minimum
+    GPGPU_Driver::getInstance().setMinFreq();
+
+    for(int i = 0; i < ELEMENTS; i++)
+    {
+        if(out[i] != in[i])
+        {
+            printk("Error at Item %d (%d != %d)!", i, out[i], in[i]);
+        }
+    }
+        
+    // free buffers
+    free((void *)in);
+    free((void *)out);
+}
+
+void run_gpgpu_test()
+{
+    // create kernel and buffer config struct
+    static kernel_config kconf;
+    static buffer_config buffconf[2];
+
+    kconf.range[0] = ELEMENTS; // number of executions
+    kconf.workgroupsize[0] = 0; // 0 = auto
+    kconf.binary = test_Gen9core_gen;
+    kconf.finish_callback = cleanUp;
+
+    // allocate buffers
+    in = (uint32_t*)uos_aligned_alloc(0x1000, kconf.range[0] * sizeof(uint32_t));
+    out = (uint32_t*)uos_aligned_alloc(0x1000, kconf.range[0] * sizeof(uint32_t));
+
+    // config buffers
+    kconf.buffCount = 2;
+    kconf.buffConfigs = buffconf;
+    kconf.buffConfigs[0].buffer = (uint32_t*)in;
+    kconf.buffConfigs[0].buffer_size = kconf.range[0] * sizeof(uint32_t);
+    kconf.buffConfigs[1].buffer = (uint32_t*)out;
+    kconf.buffConfigs[1].buffer_size = kconf.range[0] * sizeof(uint32_t);
+
+    for(int i = 0; i < ELEMENTS; i++)
+    {
+        in[i] = 0x42;
+    }
+
+    // set maximum freuqency
+    GPGPU_Driver& gpgpudriver = GPGPU_Driver::getInstance();
+    gpgpudriver.setMaxFreq();
+
+    // start gpu task
+    gpgpudriver.enqueueRun(kconf);
+}
diff --git a/repos/dde_uos-intel-gpgpu/src/gpgpu/test.h b/repos/dde_uos-intel-gpgpu/src/gpgpu/test.h
new file mode 100644
index 0000000000..f4fbc0875e
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/src/gpgpu/test.h
@@ -0,0 +1,10 @@
+#ifndef TEST_H
+#define TEST_H
+
+/**
+ * @brief run a test kernel
+ * 
+ */
+void run_gpgpu_test();
+
+#endif // TEST_H
diff --git a/repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu b/repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu
new file mode 160000
index 0000000000..eca8dac23a
--- /dev/null
+++ b/repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu
@@ -0,0 +1 @@
+Subproject commit eca8dac23a1bd6aec59c3a3142eae8a5f504bc51