mirror of
https://github.com/mmueller41/genode.git
synced 2026-01-21 20:42:56 +01:00
Merge branch 'gpgpu' into gpgpu-bench
This commit is contained in:
Marcel Lütke Dreimann
3
.gitmodules
vendored
Normal file
3
.gitmodules
vendored
Normal file
@@ -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 = https://ess.cs.uos.de/git/software/uos-intel-gpgpu.git
|
||||
0
repos/dde_uos-intel-gpgpu/README
Normal file
0
repos/dde_uos-intel-gpgpu/README
Normal file
1
repos/dde_uos-intel-gpgpu/include/gpgpu/gpgpu.h
Normal file
1
repos/dde_uos-intel-gpgpu/include/gpgpu/gpgpu.h
Normal file
@@ -0,0 +1 @@
|
||||
#include "../../src/uos-intel-gpgpu/driver/gpgpu_driver.h"
|
||||
26
repos/dde_uos-intel-gpgpu/include/gpgpu/session.h
Normal file
26
repos/dde_uos-intel-gpgpu/include/gpgpu/session.h
Normal file
@@ -0,0 +1,26 @@
|
||||
#ifndef GPGPU_SESSION
|
||||
#define GPGPU_SESSION
|
||||
|
||||
#include <session/session.h>
|
||||
#include <base/rpc.h>
|
||||
|
||||
namespace gpgpu { struct Session; }
|
||||
|
||||
struct gpgpu::Session : Genode::Session
|
||||
{
|
||||
static const char *service_name() { return "gpgpu"; }
|
||||
|
||||
enum { CAP_QUOTA = 1 };
|
||||
|
||||
virtual void say_hello() = 0;
|
||||
|
||||
/*******************
|
||||
** RPC interface **
|
||||
*******************/
|
||||
|
||||
GENODE_RPC(Rpc_say_hello, void, say_hello);
|
||||
|
||||
GENODE_RPC_INTERFACE(Rpc_say_hello);
|
||||
};
|
||||
|
||||
#endif // GPGPU_SESSION
|
||||
68
repos/dde_uos-intel-gpgpu/run/gpgpu.run
Normal file
68
repos/dde_uos-intel-gpgpu/run/gpgpu.run
Normal file
@@ -0,0 +1,68 @@
|
||||
# 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">
|
||||
<provides> <service name="gpgpu"/> </provides>
|
||||
<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
|
||||
155
repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.cc
Normal file
155
repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.cc
Normal file
@@ -0,0 +1,155 @@
|
||||
#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);
|
||||
size_t donate = size;
|
||||
ram_cap =
|
||||
retry<Out_of_ram>(
|
||||
[&] () {
|
||||
return retry<Out_of_caps>(
|
||||
[&] () { return pci.alloc_dma_buffer(size, UNCACHED); },
|
||||
[&] () { pci.upgrade_caps(2); });
|
||||
},
|
||||
[&] () {
|
||||
pci.upgrade_ram(donate);
|
||||
donate = donate * 2 > size ? 4096 : donate * 2;
|
||||
});
|
||||
mapped_base = e.rm().attach(ram_cap);
|
||||
base = pci.dma_addr(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)
|
||||
{
|
||||
return alloc.alloc_aligned(size, alignment).convert<void *>(
|
||||
|
||||
[&] (void *ptr) { return ptr; },
|
||||
|
||||
[&] (Genode::Range_allocator::Alloc_error) -> void * {
|
||||
Genode::error("[GPU] Error in driver allocation!");
|
||||
return nullptr;
|
||||
}
|
||||
);
|
||||
}
|
||||
|
||||
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(uint8_t addr)
|
||||
{
|
||||
Platform::Device_client client(dev);
|
||||
return client.config_read(addr, Platform::Device::ACCESS_32BIT);
|
||||
}
|
||||
|
||||
void gpgpu_genode::writePCI(uint8_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();
|
||||
}
|
||||
135
repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.h
Normal file
135
repos/dde_uos-intel-gpgpu/src/gpgpu/gpgpu_genode.h
Normal file
@@ -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 <legacy/x86/platform_session/connection.h>
|
||||
#include <legacy/x86/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(uint8_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(uint8_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
|
||||
91
repos/dde_uos-intel-gpgpu/src/gpgpu/main.cc
Normal file
91
repos/dde_uos-intel-gpgpu/src/gpgpu/main.cc
Normal file
@@ -0,0 +1,91 @@
|
||||
#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;
|
||||
|
||||
extern void construct_RPC(Genode::Env &env);
|
||||
|
||||
void Component::construct(Genode::Env& e)
|
||||
{
|
||||
Genode::log("Hello world: UOS Intel GPGPU!");
|
||||
Genode::log("Build: ", __TIMESTAMP__);
|
||||
|
||||
construct_RPC(e);
|
||||
return;
|
||||
|
||||
// init globals
|
||||
static gpgpu_genode gg(e);
|
||||
_global_gpgpu_genode = ≫
|
||||
|
||||
#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);
|
||||
}
|
||||
if(virt_to_phys(test) == nullptr)
|
||||
{
|
||||
Genode::error("mem phys addr NULL");
|
||||
}
|
||||
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!");
|
||||
}
|
||||
68
repos/dde_uos-intel-gpgpu/src/gpgpu/rpc.cc
Normal file
68
repos/dde_uos-intel-gpgpu/src/gpgpu/rpc.cc
Normal file
@@ -0,0 +1,68 @@
|
||||
#include <base/log.h>
|
||||
#include <base/heap.h>
|
||||
#include <root/component.h>
|
||||
#include <base/rpc_server.h>
|
||||
#include <gpgpu/session.h>
|
||||
|
||||
namespace gpgpu {
|
||||
struct Session_component;
|
||||
struct Root_component;
|
||||
struct Main;
|
||||
}
|
||||
|
||||
struct gpgpu::Session_component : Genode::Rpc_object<Session>
|
||||
{
|
||||
void say_hello() override
|
||||
{
|
||||
Genode::log("Hello from uos-intel-gpgpu!");
|
||||
}
|
||||
};
|
||||
|
||||
class gpgpu::Root_component
|
||||
:
|
||||
public Genode::Root_component<Session_component>
|
||||
{
|
||||
protected:
|
||||
Session_component *_create_session(const char *) override
|
||||
{
|
||||
return new (md_alloc()) Session_component();
|
||||
}
|
||||
|
||||
public:
|
||||
|
||||
Root_component(Genode::Entrypoint &ep,
|
||||
Genode::Allocator &alloc)
|
||||
:
|
||||
Genode::Root_component<Session_component>(ep, alloc)
|
||||
{
|
||||
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
struct gpgpu::Main
|
||||
{
|
||||
Genode::Env &env;
|
||||
|
||||
/*
|
||||
* A sliced heap is used for allocating session objects - thereby we
|
||||
* can release objects separately.
|
||||
*/
|
||||
Genode::Sliced_heap sliced_heap { env.ram(), env.rm() };
|
||||
|
||||
gpgpu::Root_component root { env.ep(), sliced_heap };
|
||||
|
||||
Main(Genode::Env &env) : env(env)
|
||||
{
|
||||
/*
|
||||
* Create a RPC object capability for the root interface and
|
||||
* announce the service to our parent.
|
||||
*/
|
||||
env.parent().announce(env.ep().manage(root));
|
||||
}
|
||||
};
|
||||
|
||||
void construct_RPC(Genode::Env &env)
|
||||
{
|
||||
static gpgpu::Main main(env);
|
||||
}
|
||||
6
repos/dde_uos-intel-gpgpu/src/gpgpu/rpc.h
Normal file
6
repos/dde_uos-intel-gpgpu/src/gpgpu/rpc.h
Normal file
@@ -0,0 +1,6 @@
|
||||
#ifndef RPC_H
|
||||
#define RPC_H
|
||||
|
||||
|
||||
|
||||
#endif // RPC_H
|
||||
67
repos/dde_uos-intel-gpgpu/src/gpgpu/stubs.cc
Normal file
67
repos/dde_uos-intel-gpgpu/src/gpgpu/stubs.cc
Normal file
@@ -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((uint8_t)addr);
|
||||
}
|
||||
|
||||
extern "C" void writePCIConfigSpace(uint32_t address, uint32_t value)
|
||||
{
|
||||
_global_gpgpu_genode->writePCI((uint8_t)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);
|
||||
}
|
||||
20
repos/dde_uos-intel-gpgpu/src/gpgpu/target.mk
Normal file
20
repos/dde_uos-intel-gpgpu/src/gpgpu/target.mk
Normal file
@@ -0,0 +1,20 @@
|
||||
TARGET = gpgpu
|
||||
REQUIRES = x86_64
|
||||
|
||||
SRC_CC = main.cc gpgpu_genode.cc stubs.cc test.cc rpc.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
|
||||
206
repos/dde_uos-intel-gpgpu/src/gpgpu/test.cc
Normal file
206
repos/dde_uos-intel-gpgpu/src/gpgpu/test.cc
Normal file
@@ -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);
|
||||
}
|
||||
10
repos/dde_uos-intel-gpgpu/src/gpgpu/test.h
Normal file
10
repos/dde_uos-intel-gpgpu/src/gpgpu/test.h
Normal file
@@ -0,0 +1,10 @@
|
||||
#ifndef TEST_H
|
||||
#define TEST_H
|
||||
|
||||
/**
|
||||
* @brief run a test kernel
|
||||
*
|
||||
*/
|
||||
void run_gpgpu_test();
|
||||
|
||||
#endif // TEST_H
|
||||
1
repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu
Submodule
1
repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu
Submodule
Submodule repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu added at 6f6cf6fb0a
0
repos/hello_gpgpu/README
Normal file
0
repos/hello_gpgpu/README
Normal file
21
repos/hello_gpgpu/include/hello_gpgpu_session/client.h
Normal file
21
repos/hello_gpgpu/include/hello_gpgpu_session/client.h
Normal file
@@ -0,0 +1,21 @@
|
||||
#ifndef HELLO_GPGPU_CLIENT_H
|
||||
#define HELLO_GPGPU_CLIENT_H
|
||||
|
||||
#include <gpgpu/session.h>
|
||||
#include <base/rpc_client.h>
|
||||
#include <base/log.h>
|
||||
|
||||
namespace gpgpu { struct Session_client; }
|
||||
|
||||
struct gpgpu::Session_client : Genode::Rpc_client<gpgpu::Session>
|
||||
{
|
||||
Session_client(Genode::Capability<gpgpu::Session> cap)
|
||||
: Genode::Rpc_client<gpgpu::Session>(cap) { }
|
||||
|
||||
void say_hello() override
|
||||
{
|
||||
call<Rpc_say_hello>();
|
||||
}
|
||||
};
|
||||
|
||||
#endif // HELLO_GPGPU_CLIENT_H
|
||||
21
repos/hello_gpgpu/include/hello_gpgpu_session/connection.h
Normal file
21
repos/hello_gpgpu/include/hello_gpgpu_session/connection.h
Normal file
@@ -0,0 +1,21 @@
|
||||
#ifndef HELLO_GPGPU_CONNECTION_H
|
||||
#define HELLO_GPGPU_CONNECTION_H
|
||||
|
||||
#include <hello_gpgpu_session/client.h>
|
||||
#include <base/connection.h>
|
||||
|
||||
namespace gpgpu { struct Connection; }
|
||||
|
||||
struct gpgpu::Connection : Genode::Connection<gpgpu::Session>, Session_client
|
||||
{
|
||||
Connection(Genode::Env &env)
|
||||
:
|
||||
/* create session */
|
||||
Genode::Connection<gpgpu::Session>(env, session(env.parent(),
|
||||
"ram_quota=6K, cap_quota=4")),
|
||||
|
||||
/* initialize RPC interface */
|
||||
Session_client(cap()) { }
|
||||
};
|
||||
|
||||
#endif // HELLO_GPGPU_CONNECTION_H
|
||||
56
repos/hello_gpgpu/run/hello_gpgpu.run
Normal file
56
repos/hello_gpgpu/run/hello_gpgpu.run
Normal file
@@ -0,0 +1,56 @@
|
||||
#
|
||||
# Build
|
||||
#
|
||||
|
||||
build { core init gpgpu timer hello_gpgpu }
|
||||
|
||||
create_boot_directory
|
||||
|
||||
#
|
||||
# Generate config
|
||||
#
|
||||
|
||||
install_config {
|
||||
<config>
|
||||
<parent-provides>
|
||||
<service name="LOG"/>
|
||||
<service name="PD"/>
|
||||
<service name="CPU"/>
|
||||
<service name="ROM"/>
|
||||
<service name="RAM"/>
|
||||
<service name="CAP"/>
|
||||
<service name="RM"/>
|
||||
</parent-provides>
|
||||
<default-route>
|
||||
<any-service> <parent/> <any-child/> </any-service>
|
||||
</default-route>
|
||||
<default caps="200"/>
|
||||
<start name="gpgpu">
|
||||
<resource name="RAM" quantum="10M"/>
|
||||
<provides><service name="gpgpu"/></provides>
|
||||
</start>
|
||||
<start name="timer">
|
||||
<resource name="RAM" quantum="1M"/>
|
||||
<provides><service name="Timer"/></provides>
|
||||
<route>
|
||||
<any-service><parent/><any-child/></any-service>
|
||||
</route>
|
||||
</start>
|
||||
<start name="hello_gpgpu">
|
||||
<resource name="RAM" quantum="1024M"/>
|
||||
<config>
|
||||
<vfs> <dir name="dev"> <log/> <inline name="rtc">2022-07-20 14:30</inline> </dir> </vfs>
|
||||
<libc stdout="/dev/log" stderr="/dev/log" rtc="/dev/rtc"/>
|
||||
</config>
|
||||
</start>
|
||||
</config>}
|
||||
|
||||
#
|
||||
# Boot image
|
||||
#
|
||||
|
||||
build_boot_image { core ld.lib.so libc.lib.so vfs.lib.so init gpgpu timer hello_gpgpu }
|
||||
|
||||
append qemu_args " -nographic -m 4G"
|
||||
|
||||
run_genode_until "hello gpgpu completed.*\n" 15
|
||||
418
repos/hello_gpgpu/src/hello_gpgpu/2mm.cc
Normal file
418
repos/hello_gpgpu/src/hello_gpgpu/2mm.cc
Normal file
@@ -0,0 +1,418 @@
|
||||
/**
|
||||
* 2mm.c: This file is part of the PolyBench/GPU 1.0 test suite.
|
||||
*
|
||||
*
|
||||
* Contact: Scott Grauer-Gray <sgrauerg@gmail.com>
|
||||
* Will Killian <killian@udel.edu>
|
||||
* Louis-Noel Pouchet <pouchet@cse.ohio-state.edu>
|
||||
* Web address: http://www.cse.ohio-state.edu/~pouchet/software/polybench/GPU
|
||||
*/
|
||||
|
||||
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <time.h>
|
||||
#include <sys/time.h>
|
||||
#include <math.h>
|
||||
|
||||
#define CL_TARGET_OPENCL_VERSION 100
|
||||
#ifdef __APPLE__
|
||||
#include <OpenCL/opencl.h>
|
||||
#else
|
||||
#include "CL/cl.h"
|
||||
#endif
|
||||
|
||||
#define POLYBENCH_TIME 1
|
||||
|
||||
//select the OpenCL device to use (can be GPU, CPU, or Accelerator such as Intel Xeon Phi)
|
||||
#define OPENCL_DEVICE_SELECTION CL_DEVICE_TYPE_GPU
|
||||
|
||||
#include "2mm.h"
|
||||
#include "polybench.h"
|
||||
#include "polybenchUtilFuncts.h"
|
||||
|
||||
//define the error threshold for the results "not matching"
|
||||
#define PERCENT_DIFF_ERROR_THRESHOLD 1.05
|
||||
|
||||
#define MAX_SOURCE_SIZE (0x100000)
|
||||
|
||||
|
||||
#if defined(cl_khr_fp64) // Khronos extension available?
|
||||
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
|
||||
#elif defined(cl_amd_fp64) // AMD extension available?
|
||||
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
|
||||
#endif
|
||||
|
||||
#include "2mm_kernel.h"
|
||||
|
||||
char str_temp[1024];
|
||||
|
||||
cl_platform_id platform_id;
|
||||
cl_device_id device_id;
|
||||
cl_uint num_devices;
|
||||
cl_uint num_platforms;
|
||||
cl_int errcode;
|
||||
cl_context clGPUContext;
|
||||
cl_kernel clKernel1;
|
||||
cl_kernel clKernel2;
|
||||
cl_command_queue clCommandQue;
|
||||
cl_program clProgram;
|
||||
cl_mem tmp_mem_obj;
|
||||
cl_mem a_mem_obj;
|
||||
cl_mem b_mem_obj;
|
||||
cl_mem c_mem_obj;
|
||||
cl_mem dOutputFromGpu_mem_obj;
|
||||
|
||||
FILE *fp;
|
||||
char *source_str;
|
||||
size_t source_size;
|
||||
|
||||
#define RUN_ON_CPU
|
||||
|
||||
|
||||
void compareResults(int ni, int nl, DATA_TYPE POLYBENCH_2D(D, NI, NL, ni, nl), DATA_TYPE POLYBENCH_2D(D_outputFromGpu, NI, NL, ni, nl))
|
||||
{
|
||||
int i,j,fail;
|
||||
fail = 0;
|
||||
|
||||
for (i=0; i < ni; i++)
|
||||
{
|
||||
for (j=0; j < nl; j++)
|
||||
{
|
||||
if (percentDiff(D[i][j], D_outputFromGpu[i][j]) > PERCENT_DIFF_ERROR_THRESHOLD)
|
||||
{
|
||||
fail++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// print results
|
||||
printf("Non-Matching CPU-GPU Outputs Beyond Error Threshold of %4.2f Percent: %d\n", PERCENT_DIFF_ERROR_THRESHOLD, fail);
|
||||
}
|
||||
|
||||
|
||||
void read_cl_file()
|
||||
{
|
||||
// Load the kernel source code into the array source_str
|
||||
/*fp = fopen("2mm.cl", "r");
|
||||
if (!fp) {
|
||||
fprintf(stderr, "Failed to load kernel.\n");
|
||||
exit(1);
|
||||
}
|
||||
source_str = (char*)malloc(MAX_SOURCE_SIZE);
|
||||
source_size = fread( source_str, 1, MAX_SOURCE_SIZE, fp);
|
||||
fclose( fp );*/
|
||||
}
|
||||
|
||||
|
||||
void init_array(int ni, int nj, int nk, int nl, DATA_TYPE *alpha, DATA_TYPE *beta, DATA_TYPE POLYBENCH_2D(A, NI, NK, ni, nk),
|
||||
DATA_TYPE POLYBENCH_2D(B, NK, NJ, nk, nj), DATA_TYPE POLYBENCH_2D(C, NL, NJ, nl, nj),
|
||||
DATA_TYPE POLYBENCH_2D(D, NI, NL, ni, nl), DATA_TYPE POLYBENCH_2D(Dgpu, NI, NL, ni, nl))
|
||||
{
|
||||
int i, j;
|
||||
|
||||
*alpha = 32412;
|
||||
*beta = 2123;
|
||||
|
||||
for (i = 0; i < ni; i++)
|
||||
{
|
||||
for (j = 0; j < nk; j++)
|
||||
{
|
||||
A[i][j] = ((DATA_TYPE) (i*j)) / NI;
|
||||
}
|
||||
}
|
||||
|
||||
for (i = 0; i < nk; i++)
|
||||
{
|
||||
for (j = 0; j < nj; j++)
|
||||
{
|
||||
B[i][j] = ((DATA_TYPE) (i*(j+1))) / NJ;
|
||||
}
|
||||
}
|
||||
|
||||
for (i = 0; i < nl; i++)
|
||||
{
|
||||
for (j = 0; j < nj; j++)
|
||||
{
|
||||
C[i][j] = ((DATA_TYPE) (i*(j+3))) / NL;
|
||||
}
|
||||
}
|
||||
|
||||
for (i = 0; i < ni; i++)
|
||||
{
|
||||
for (j = 0; j < nl; j++)
|
||||
{
|
||||
D[i][j] = ((DATA_TYPE) (i*(j+2))) / NK;
|
||||
Dgpu[i][j] = ((DATA_TYPE) (i*(j+2))) / NK;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void cl_initialization()
|
||||
{
|
||||
// Get platform and device information
|
||||
errcode = clGetPlatformIDs(1, &platform_id, &num_platforms);
|
||||
if(errcode == CL_SUCCESS) printf("number of platforms is %d\n",num_platforms);
|
||||
else printf("Error getting platform IDs\n");
|
||||
|
||||
errcode = clGetPlatformInfo(platform_id,CL_PLATFORM_NAME, sizeof(str_temp), str_temp,NULL);
|
||||
if(errcode == CL_SUCCESS) printf("platform name is %s\n",str_temp);
|
||||
else printf("Error getting platform name\n");
|
||||
|
||||
errcode = clGetPlatformInfo(platform_id, CL_PLATFORM_VERSION, sizeof(str_temp), str_temp,NULL);
|
||||
if(errcode == CL_SUCCESS) printf("platform version is %s\n",str_temp);
|
||||
else printf("Error getting platform version\n");
|
||||
|
||||
errcode = clGetDeviceIDs( platform_id, OPENCL_DEVICE_SELECTION, 1, &device_id, &num_devices);
|
||||
if(errcode == CL_SUCCESS) printf("number of devices is %d\n", num_devices);
|
||||
else printf("Error getting device IDs\n");
|
||||
|
||||
errcode = clGetDeviceInfo(device_id,CL_DEVICE_NAME, sizeof(str_temp), str_temp,NULL);
|
||||
if(errcode == CL_SUCCESS) printf("device name is %s\n",str_temp);
|
||||
else printf("Error getting device name\n");
|
||||
|
||||
// Create an OpenCL context
|
||||
clGPUContext = clCreateContext( NULL, 1, &device_id, NULL, NULL, &errcode);
|
||||
if(errcode != CL_SUCCESS) printf("Error in creating context\n");
|
||||
|
||||
//Create a command-queue
|
||||
clCommandQue = clCreateCommandQueue(clGPUContext, device_id, 0, &errcode);
|
||||
if(errcode != CL_SUCCESS) printf("Error in creating command queue\n");
|
||||
}
|
||||
|
||||
|
||||
void cl_mem_init(DATA_TYPE POLYBENCH_2D(tmp, NI,NJ,ni,nj), DATA_TYPE POLYBENCH_2D(A, NI,NK,ni,nk), DATA_TYPE POLYBENCH_2D(B, NK,NJ,nk,nj),
|
||||
DATA_TYPE POLYBENCH_2D(C, NL,NJ,nl,nj), DATA_TYPE POLYBENCH_2D(D_outputFromGpu,NI,NL,ni,nl))
|
||||
{
|
||||
tmp_mem_obj = clCreateBuffer(clGPUContext, CL_MEM_READ_WRITE, sizeof(DATA_TYPE) * NI * NJ, NULL, &errcode);
|
||||
a_mem_obj = clCreateBuffer(clGPUContext, CL_MEM_READ_ONLY, sizeof(DATA_TYPE) * NI * NK, NULL, &errcode);
|
||||
b_mem_obj = clCreateBuffer(clGPUContext, CL_MEM_READ_ONLY, sizeof(DATA_TYPE) * NK * NJ, NULL, &errcode);
|
||||
c_mem_obj = clCreateBuffer(clGPUContext, CL_MEM_READ_WRITE, sizeof(DATA_TYPE) * NL * NJ, NULL, &errcode);
|
||||
dOutputFromGpu_mem_obj = clCreateBuffer(clGPUContext, CL_MEM_READ_WRITE, sizeof(DATA_TYPE) * NI * NL, NULL, &errcode);
|
||||
if(errcode != CL_SUCCESS) printf("Error in creating buffers\n");
|
||||
|
||||
errcode = clEnqueueWriteBuffer(clCommandQue, tmp_mem_obj, CL_TRUE, 0, sizeof(DATA_TYPE) * NI * NJ, tmp, 0, NULL, NULL);
|
||||
errcode = clEnqueueWriteBuffer(clCommandQue, a_mem_obj, CL_TRUE, 0, sizeof(DATA_TYPE) * NI * NK, A, 0, NULL, NULL);
|
||||
errcode = clEnqueueWriteBuffer(clCommandQue, b_mem_obj, CL_TRUE, 0, sizeof(DATA_TYPE) * NK * NJ, B, 0, NULL, NULL);
|
||||
errcode = clEnqueueWriteBuffer(clCommandQue, c_mem_obj, CL_TRUE, 0, sizeof(DATA_TYPE) * NL * NJ, C, 0, NULL, NULL);
|
||||
errcode = clEnqueueWriteBuffer(clCommandQue, dOutputFromGpu_mem_obj, CL_TRUE, 0, sizeof(DATA_TYPE) * NI * NL, D_outputFromGpu, 0, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS)printf("Error in writing buffers\n");
|
||||
}
|
||||
|
||||
|
||||
void cl_load_prog()
|
||||
{
|
||||
// Create a program from the kernel source
|
||||
const size_t kernel_size = __2mm_Gen9core_gen_len;
|
||||
const unsigned char* kernel_bin = __2mm_Gen9core_gen;
|
||||
clProgram = clCreateProgramWithBinary(clGPUContext, 1, &device_id, &kernel_size, &kernel_bin, NULL, &errcode);
|
||||
//clProgram = clCreateProgramWithSource(clGPUContext, 1, (const char **)&source_str, (const size_t *)&source_size, &errcode);
|
||||
|
||||
if(errcode != CL_SUCCESS) printf("Error in creating program\n");
|
||||
|
||||
// Build the program
|
||||
errcode = clBuildProgram(clProgram, 1, &device_id, NULL, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS) printf("Error in building program\n");
|
||||
|
||||
// Create the OpenCL kernel
|
||||
clKernel1 = clCreateKernel(clProgram, "mm2_kernel1", &errcode);
|
||||
if(errcode != CL_SUCCESS) printf("Error in creating kernel\n");
|
||||
|
||||
clKernel2 = clCreateKernel(clProgram, "mm2_kernel2", &errcode);
|
||||
if(errcode != CL_SUCCESS) printf("Error in creating kernel\n");
|
||||
clFinish(clCommandQue);
|
||||
}
|
||||
|
||||
|
||||
void cl_launch_kernel(int ni, int nj, int nk, int nl, DATA_TYPE alpha, DATA_TYPE beta)
|
||||
{
|
||||
size_t localWorkSize[2], globalWorkSize[2];
|
||||
localWorkSize[0] = DIM_LOCAL_WORK_GROUP_X;
|
||||
localWorkSize[1] = DIM_LOCAL_WORK_GROUP_Y;
|
||||
globalWorkSize[0] = (size_t)ceil(((float)NI) / ((float)DIM_LOCAL_WORK_GROUP_X)) * DIM_LOCAL_WORK_GROUP_X;
|
||||
globalWorkSize[1] = (size_t)ceil(((float)NL) / ((float)DIM_LOCAL_WORK_GROUP_Y)) * DIM_LOCAL_WORK_GROUP_Y;
|
||||
|
||||
/* Start timer. */
|
||||
polybench_start_instruments;
|
||||
|
||||
// Set the arguments of the kernel
|
||||
errcode = clSetKernelArg(clKernel1, 0, sizeof(cl_mem), (void *)&tmp_mem_obj);
|
||||
errcode |= clSetKernelArg(clKernel1, 1, sizeof(cl_mem), (void *)&a_mem_obj);
|
||||
errcode |= clSetKernelArg(clKernel1, 2, sizeof(cl_mem), (void *)&b_mem_obj);
|
||||
errcode |= clSetKernelArg(clKernel1, 3, sizeof(int), (void *)&ni);
|
||||
errcode |= clSetKernelArg(clKernel1, 4, sizeof(int), (void *)&nj);
|
||||
errcode |= clSetKernelArg(clKernel1, 5, sizeof(int), (void *)&nk);
|
||||
errcode |= clSetKernelArg(clKernel1, 6, sizeof(int), (void *)&nl);
|
||||
errcode |= clSetKernelArg(clKernel1, 7, sizeof(DATA_TYPE), (void *)&alpha);
|
||||
errcode |= clSetKernelArg(clKernel1, 8, sizeof(DATA_TYPE), (void *)&beta);
|
||||
if(errcode != CL_SUCCESS) printf("Error in seting arguments\n");
|
||||
// Execute the OpenCL kernel
|
||||
errcode = clEnqueueNDRangeKernel(clCommandQue, clKernel1, 2, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS) printf("Error in launching kernel\n");
|
||||
clEnqueueBarrier(clCommandQue);
|
||||
|
||||
globalWorkSize[0] = (size_t)ceil(((float)NI) / ((float)DIM_LOCAL_WORK_GROUP_X)) * DIM_LOCAL_WORK_GROUP_X;
|
||||
globalWorkSize[1] = (size_t)ceil(((float)NL) / ((float)DIM_LOCAL_WORK_GROUP_Y)) * DIM_LOCAL_WORK_GROUP_Y;
|
||||
|
||||
errcode = clSetKernelArg(clKernel2, 0, sizeof(cl_mem), (void *)&tmp_mem_obj);
|
||||
errcode |= clSetKernelArg(clKernel2, 1, sizeof(cl_mem), (void *)&c_mem_obj);
|
||||
errcode |= clSetKernelArg(clKernel2, 2, sizeof(cl_mem), (void *)&dOutputFromGpu_mem_obj);
|
||||
errcode |= clSetKernelArg(clKernel2, 3, sizeof(int), (void *)&ni);
|
||||
errcode |= clSetKernelArg(clKernel2, 4, sizeof(int), (void *)&nj);
|
||||
errcode |= clSetKernelArg(clKernel2, 5, sizeof(int), (void *)&nk);
|
||||
errcode |= clSetKernelArg(clKernel2, 6, sizeof(int), (void *)&nl);
|
||||
errcode |= clSetKernelArg(clKernel2, 7, sizeof(DATA_TYPE), (void *)&alpha);
|
||||
errcode |= clSetKernelArg(clKernel2, 8, sizeof(DATA_TYPE), (void *)&beta);
|
||||
if(errcode != CL_SUCCESS) printf("Error in seting arguments\n");
|
||||
|
||||
// Execute the OpenCL kernel
|
||||
errcode = clEnqueueNDRangeKernel(clCommandQue, clKernel2, 2, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS) printf("Error in launching kernel\n");
|
||||
clFinish(clCommandQue);
|
||||
|
||||
/* Stop and print timer. */
|
||||
printf("GPU Time in seconds:\n");
|
||||
polybench_stop_instruments;
|
||||
polybench_print_instruments;
|
||||
}
|
||||
|
||||
|
||||
void cl_clean_up()
|
||||
{
|
||||
// Clean up
|
||||
errcode = clFlush(clCommandQue);
|
||||
errcode = clFinish(clCommandQue);
|
||||
errcode = clReleaseKernel(clKernel1);
|
||||
errcode = clReleaseKernel(clKernel2);
|
||||
errcode = clReleaseProgram(clProgram);
|
||||
errcode = clReleaseMemObject(tmp_mem_obj);
|
||||
errcode = clReleaseMemObject(a_mem_obj);
|
||||
errcode = clReleaseMemObject(b_mem_obj);
|
||||
errcode = clReleaseMemObject(c_mem_obj);
|
||||
errcode = clReleaseMemObject(dOutputFromGpu_mem_obj);
|
||||
errcode = clReleaseCommandQueue(clCommandQue);
|
||||
errcode = clReleaseContext(clGPUContext);
|
||||
if(errcode != CL_SUCCESS) printf("Error in cleanup\n");
|
||||
}
|
||||
|
||||
|
||||
void mm2_cpu(int ni, int nj, int nk, int nl,
|
||||
DATA_TYPE alpha,
|
||||
DATA_TYPE beta,
|
||||
DATA_TYPE POLYBENCH_2D(tmp,NI,NJ,ni,nj),
|
||||
DATA_TYPE POLYBENCH_2D(A,NI,NK,ni,nk),
|
||||
DATA_TYPE POLYBENCH_2D(B,NK,NJ,nk,nj),
|
||||
DATA_TYPE POLYBENCH_2D(C,NL,NJ,nl,nj),
|
||||
DATA_TYPE POLYBENCH_2D(D,NI,NL,ni,nl))
|
||||
{
|
||||
int i, j, k;
|
||||
|
||||
/* D := alpha*A*B*C + beta*D */
|
||||
for (i = 0; i < _PB_NI; i++)
|
||||
{
|
||||
for (j = 0; j < _PB_NJ; j++)
|
||||
{
|
||||
tmp[i][j] = 0;
|
||||
for (k = 0; k < _PB_NK; ++k)
|
||||
{
|
||||
tmp[i][j] += alpha * A[i][k] * B[k][j];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
for (i = 0; i < _PB_NI; i++)
|
||||
{
|
||||
for (j = 0; j < _PB_NL; j++)
|
||||
{
|
||||
D[i][j] *= beta;
|
||||
for (k = 0; k < _PB_NJ; ++k)
|
||||
{
|
||||
D[i][j] += tmp[i][k] * C[k][j];
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* DCE code. Must scan the entire live-out data.
|
||||
Can be used also to check the correctness of the output. */
|
||||
static
|
||||
void print_array(int ni, int nl,
|
||||
DATA_TYPE POLYBENCH_2D(D,NI,NL,ni,nl))
|
||||
{
|
||||
int i, j;
|
||||
|
||||
for (i = 0; i < ni; i++)
|
||||
for (j = 0; j < nl; j++) {
|
||||
fprintf (stderr, DATA_PRINTF_MODIFIER, D[i][j]);
|
||||
if ((i * ni + j) % 20 == 0) fprintf (stderr, "\n");
|
||||
}
|
||||
fprintf (stderr, "\n");
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
/* Retrieve problem size. */
|
||||
int ni = NI;
|
||||
int nj = NJ;
|
||||
int nk = NK;
|
||||
int nl = NL;
|
||||
|
||||
/* Variable declaration/allocation. */
|
||||
DATA_TYPE alpha;
|
||||
DATA_TYPE beta;
|
||||
POLYBENCH_2D_ARRAY_DECL(tmp,DATA_TYPE,NI,NJ,ni,nj);
|
||||
POLYBENCH_2D_ARRAY_DECL(A,DATA_TYPE,NI,NK,ni,nk);
|
||||
POLYBENCH_2D_ARRAY_DECL(B,DATA_TYPE,NK,NJ,nk,nj);
|
||||
POLYBENCH_2D_ARRAY_DECL(C,DATA_TYPE,NL,NJ,nl,nj);
|
||||
POLYBENCH_2D_ARRAY_DECL(D,DATA_TYPE,NI,NL,ni,nl);
|
||||
POLYBENCH_2D_ARRAY_DECL(D_outputFromGpu,DATA_TYPE,NI,NL,ni,nl);
|
||||
|
||||
/* Initialize array(s). */
|
||||
init_array(ni, nj, nk, nl, &alpha, &beta, POLYBENCH_ARRAY(A), POLYBENCH_ARRAY(B), POLYBENCH_ARRAY(C), POLYBENCH_ARRAY(D), POLYBENCH_ARRAY(D_outputFromGpu));
|
||||
|
||||
read_cl_file();
|
||||
cl_initialization();
|
||||
cl_mem_init(POLYBENCH_ARRAY(tmp), POLYBENCH_ARRAY(A), POLYBENCH_ARRAY(B), POLYBENCH_ARRAY(C), POLYBENCH_ARRAY(D_outputFromGpu));
|
||||
cl_load_prog();
|
||||
|
||||
cl_launch_kernel(ni, nj, nk, nl, alpha, beta);
|
||||
|
||||
errcode = clEnqueueReadBuffer(clCommandQue, dOutputFromGpu_mem_obj, CL_TRUE, 0, sizeof(DATA_TYPE) * NI * NL, POLYBENCH_ARRAY(D_outputFromGpu), 0, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS) printf("Error in reading GPU mem\n");
|
||||
|
||||
|
||||
#ifdef RUN_ON_CPU
|
||||
|
||||
/* Start timer. */
|
||||
polybench_start_instruments;
|
||||
|
||||
mm2_cpu(ni, nj, nk, nl, alpha, beta, POLYBENCH_ARRAY(tmp), POLYBENCH_ARRAY(A), POLYBENCH_ARRAY(B), POLYBENCH_ARRAY(C), POLYBENCH_ARRAY(D));
|
||||
|
||||
/* Stop and print timer. */
|
||||
printf("CPU Time in seconds:\n");
|
||||
polybench_stop_instruments;
|
||||
polybench_print_instruments;
|
||||
|
||||
compareResults(ni, nl, POLYBENCH_ARRAY(D), POLYBENCH_ARRAY(D_outputFromGpu));
|
||||
|
||||
#else //prevent dead code elimination
|
||||
|
||||
polybench_prevent_dce(print_array(ni, nl, POLYBENCH_ARRAY(D_outputFromGpu)));
|
||||
|
||||
#endif //RUN_ON_CPU
|
||||
|
||||
cl_clean_up();
|
||||
|
||||
POLYBENCH_FREE_ARRAY(tmp);
|
||||
POLYBENCH_FREE_ARRAY(A);
|
||||
POLYBENCH_FREE_ARRAY(B);
|
||||
POLYBENCH_FREE_ARRAY(C);
|
||||
POLYBENCH_FREE_ARRAY(D);
|
||||
POLYBENCH_FREE_ARRAY(D_outputFromGpu);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
#include "polybench.cc"
|
||||
73
repos/hello_gpgpu/src/hello_gpgpu/2mm.h
Normal file
73
repos/hello_gpgpu/src/hello_gpgpu/2mm.h
Normal file
@@ -0,0 +1,73 @@
|
||||
/**
|
||||
* 2mm.h: This file is part of the PolyBench/GPU 1.0 test suite.
|
||||
*
|
||||
*
|
||||
* Contact: Scott Grauer-Gray <sgrauerg@gmail.com>
|
||||
* Will Killian <killian@udel.edu>
|
||||
* Louis-Noel Pouchet <pouchet@cse.ohio-state.edu>
|
||||
* Web address: http://www.cse.ohio-state.edu/~pouchet/software/polybench/GPU
|
||||
*/
|
||||
|
||||
#ifndef TWOMM_H
|
||||
# define TWOMM_H
|
||||
|
||||
/* Default to STANDARD_DATASET. */
|
||||
# if !defined(MINI_DATASET) && !defined(SMALL_DATASET) && !defined(LARGE_DATASET) && !defined(EXTRALARGE_DATASET)
|
||||
# define SMALL_DATASET
|
||||
# endif
|
||||
|
||||
/* Do not define anything if the user manually defines the size. */
|
||||
# if !defined(NI) && !defined(NJ) && !defined(NK) && !defined(NL)
|
||||
/* Define the possible dataset sizes. */
|
||||
# ifdef MINI_DATASET
|
||||
# define NI 256
|
||||
# define NJ 256
|
||||
# define NK 256
|
||||
# define NL 256
|
||||
# endif
|
||||
|
||||
# ifdef SMALL_DATASET
|
||||
# define NI 512
|
||||
# define NJ 512
|
||||
# define NK 512
|
||||
# define NL 512
|
||||
# endif
|
||||
|
||||
# ifdef STANDARD_DATASET /* Default if unspecified. */
|
||||
# define NI 1024
|
||||
# define NJ 1024
|
||||
# define NK 1024
|
||||
# define NL 1024
|
||||
# endif
|
||||
|
||||
# ifdef LARGE_DATASET
|
||||
# define NI 2048
|
||||
# define NJ 2048
|
||||
# define NK 2048
|
||||
# define NL 2048
|
||||
# endif
|
||||
|
||||
# ifdef EXTRALARGE_DATASET
|
||||
# define NI 4096
|
||||
# define NJ 4096
|
||||
# define NK 4096
|
||||
# define NL 4096
|
||||
# endif
|
||||
# endif /* !N */
|
||||
|
||||
# define _PB_NI POLYBENCH_LOOP_BOUND(NI,ni)
|
||||
# define _PB_NJ POLYBENCH_LOOP_BOUND(NJ,nj)
|
||||
# define _PB_NK POLYBENCH_LOOP_BOUND(NK,nk)
|
||||
# define _PB_NL POLYBENCH_LOOP_BOUND(NL,nl)
|
||||
|
||||
# ifndef DATA_TYPE
|
||||
# define DATA_TYPE float
|
||||
# define DATA_PRINTF_MODIFIER "%0.2lf "
|
||||
# endif
|
||||
|
||||
/* Thread block dimensions */
|
||||
#define DIM_LOCAL_WORK_GROUP_X 32
|
||||
#define DIM_LOCAL_WORK_GROUP_Y 8
|
||||
|
||||
|
||||
#endif /* !TWOMM*/
|
||||
570
repos/hello_gpgpu/src/hello_gpgpu/2mm_kernel.h
Normal file
570
repos/hello_gpgpu/src/hello_gpgpu/2mm_kernel.h
Normal file
@@ -0,0 +1,570 @@
|
||||
unsigned char __2mm_Gen9core_gen[] = {
|
||||
0x43, 0x54, 0x4e, 0x49, 0x39, 0x04, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x88, 0x34, 0xad, 0xe9, 0x49, 0x73, 0x32, 0xcf,
|
||||
0xd7, 0x42, 0x73, 0xa9, 0x0c, 0x00, 0x00, 0x00, 0x9c, 0x07, 0x00, 0x00,
|
||||
0x80, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00,
|
||||
0xcc, 0x00, 0x00, 0x00, 0xf0, 0x03, 0x00, 0x00, 0x6d, 0x6d, 0x32, 0x5f,
|
||||
0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x31, 0x00, 0x01, 0x00, 0x60, 0x00,
|
||||
0x0c, 0x02, 0xa0, 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, 0x80, 0x2d, 0x20, 0x00, 0x7e, 0x0a, 0x05,
|
||||
0x41, 0x00, 0x00, 0x00, 0x2c, 0x0a, 0xc0, 0x20, 0x44, 0x01, 0x00, 0x0a,
|
||||
0xb8, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x4c, 0x16, 0xc4, 0x20,
|
||||
0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x01, 0x00, 0x40, 0x20, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x80, 0x2f, 0xc0, 0x0f, 0x00, 0x12, 0x40, 0x00, 0xb1, 0x00,
|
||||
0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0xe0, 0x21, 0xc0, 0x00, 0x00, 0x12,
|
||||
0x80, 0x00, 0xb1, 0x00, 0x40, 0x00, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x21,
|
||||
0xc0, 0x0f, 0x00, 0x12, 0x20, 0x00, 0xb1, 0x00, 0x40, 0x00, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x00, 0x2f, 0xc0, 0x00, 0x00, 0x12, 0x60, 0x00, 0xb1, 0x00,
|
||||
0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x40, 0x2f, 0x80, 0x0f, 0x8d, 0x0a,
|
||||
0xe0, 0x00, 0x00, 0x00, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x20, 0x22,
|
||||
0xe0, 0x01, 0x8d, 0x0a, 0xe4, 0x00, 0x00, 0x00, 0x40, 0x96, 0x01, 0x20,
|
||||
0x07, 0x0d, 0x0b, 0x07, 0x40, 0x96, 0x2d, 0x20, 0x07, 0x76, 0x78, 0x07,
|
||||
0x10, 0x20, 0x80, 0x05, 0x20, 0x0a, 0x00, 0x20, 0x40, 0x0f, 0x8d, 0x0a,
|
||||
0x1c, 0x01, 0x00, 0x00, 0x10, 0x20, 0x80, 0x05, 0x22, 0x0a, 0x00, 0x20,
|
||||
0x20, 0x02, 0x8d, 0x0a, 0x18, 0x01, 0x00, 0x00, 0x10, 0x00, 0x80, 0x05,
|
||||
0x22, 0x0a, 0x00, 0x20, 0xa0, 0x01, 0x8d, 0x0a, 0x1c, 0x01, 0x00, 0x00,
|
||||
0x02, 0x20, 0x81, 0x00, 0x48, 0x12, 0xe0, 0x2c, 0xc4, 0x00, 0x00, 0x16,
|
||||
0x00, 0x00, 0x00, 0x00, 0x02, 0x20, 0x81, 0x00, 0x4a, 0x12, 0xe0, 0x23,
|
||||
0xc4, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x81, 0x05,
|
||||
0x22, 0x0a, 0x00, 0x20, 0xc0, 0x0e, 0x8d, 0x0a, 0x18, 0x01, 0x00, 0x00,
|
||||
0x05, 0x20, 0x80, 0x02, 0x42, 0x12, 0x00, 0x20, 0xe0, 0x03, 0xb1, 0x12,
|
||||
0xe0, 0x0c, 0xb1, 0x00, 0x22, 0x00, 0xa1, 0x00, 0x02, 0x0e, 0x00, 0x20,
|
||||
0xa0, 0x02, 0x00, 0x00, 0xa0, 0x02, 0x00, 0x00, 0x41, 0x96, 0x79, 0x20,
|
||||
0x07, 0x74, 0x76, 0x08, 0x41, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x22,
|
||||
0x20, 0x02, 0x8d, 0x0a, 0x1c, 0x01, 0x00, 0x00, 0x01, 0x00, 0x80, 0x00,
|
||||
0xe8, 0x3e, 0x40, 0x2e, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0x01, 0x20, 0x80, 0x00, 0xe8, 0x3e, 0xa0, 0x22, 0x00, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x40, 0x96, 0x01, 0x20, 0xe7, 0x74, 0x74, 0x0d,
|
||||
0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x22, 0x60, 0x02, 0x8d, 0x0a,
|
||||
0x40, 0x0f, 0x8d, 0x00, 0x10, 0x00, 0x80, 0x03, 0x24, 0x0a, 0x00, 0x20,
|
||||
0x20, 0x01, 0x00, 0x1e, 0x00, 0x00, 0x00, 0x00, 0x10, 0x20, 0x80, 0x03,
|
||||
0x24, 0x0a, 0x00, 0x20, 0x20, 0x01, 0x00, 0x1e, 0x00, 0x00, 0x00, 0x00,
|
||||
0x09, 0x00, 0x80, 0x00, 0x28, 0x0a, 0x80, 0x2e, 0x80, 0x0e, 0x8d, 0x1e,
|
||||
0x02, 0x00, 0x02, 0x00, 0x09, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x22,
|
||||
0x60, 0x02, 0x8d, 0x1e, 0x02, 0x00, 0x02, 0x00, 0x40, 0x96, 0x6d, 0x20,
|
||||
0x07, 0x74, 0x74, 0x09, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x22,
|
||||
0x60, 0x02, 0x8d, 0x0a, 0x30, 0x01, 0x00, 0x00, 0x33, 0x00, 0x80, 0x0c,
|
||||
0x70, 0x20, 0x07, 0x00, 0x82, 0x0e, 0x00, 0x00, 0x00, 0x5e, 0x02, 0x04,
|
||||
0x33, 0x20, 0x80, 0x0c, 0x70, 0x50, 0x01, 0x00, 0x62, 0x02, 0x00, 0x00,
|
||||
0x00, 0x5e, 0x02, 0x04, 0x20, 0x00, 0x11, 0x00, 0x04, 0x00, 0x00, 0x34,
|
||||
0x00, 0x14, 0x00, 0x0e, 0xb8, 0x01, 0x00, 0x00, 0x41, 0x96, 0x01, 0x20,
|
||||
0x07, 0x70, 0x76, 0x09, 0x41, 0x20, 0x80, 0x00, 0x28, 0x0a, 0xe0, 0x22,
|
||||
0x20, 0x02, 0x8d, 0x0a, 0x20, 0x01, 0x00, 0x00, 0x01, 0x00, 0x80, 0x00,
|
||||
0xe8, 0x3e, 0xc0, 0x2d, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0x01, 0x20, 0x80, 0x00, 0xe8, 0x3e, 0x20, 0x23, 0x00, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x2c, 0x1e, 0xc4, 0x2f,
|
||||
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x40, 0x96, 0x2d, 0x20,
|
||||
0x07, 0x1b, 0x70, 0x7e, 0x41, 0x00, 0x00, 0x00, 0x2c, 0x0a, 0xc8, 0x2f,
|
||||
0xc4, 0x0f, 0x00, 0x0a, 0x1c, 0x01, 0x00, 0x00, 0x40, 0x20, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x80, 0x2d, 0xe0, 0x02, 0x8d, 0x0a, 0xc4, 0x0f, 0x00, 0x00,
|
||||
0x40, 0x00, 0x00, 0x00, 0x2c, 0x0a, 0xc4, 0x2f, 0xc4, 0x0f, 0x00, 0x1e,
|
||||
0x01, 0x00, 0x01, 0x00, 0x09, 0x00, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x23,
|
||||
0x60, 0x03, 0x8d, 0x1e, 0x02, 0x00, 0x02, 0x00, 0x40, 0x96, 0x19, 0x20,
|
||||
0xe0, 0x24, 0x7e, 0x0d, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x00, 0x2d,
|
||||
0xc8, 0x0f, 0x00, 0x0a, 0x40, 0x0f, 0x8d, 0x00, 0x09, 0x20, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x80, 0x2d, 0x80, 0x0d, 0x8d, 0x1e, 0x02, 0x00, 0x02, 0x00,
|
||||
0x40, 0x96, 0x71, 0x20, 0x07, 0x1b, 0x1b, 0x09, 0x09, 0x00, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x80, 0x24, 0x80, 0x04, 0x8d, 0x1e, 0x02, 0x00, 0x02, 0x00,
|
||||
0x09, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x00, 0x2d, 0x00, 0x0d, 0x8d, 0x1e,
|
||||
0x02, 0x00, 0x02, 0x00, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x80, 0x2d,
|
||||
0x80, 0x0d, 0x8d, 0x0a, 0x34, 0x01, 0x00, 0x00, 0x31, 0x00, 0x80, 0x0c,
|
||||
0x68, 0x02, 0xa0, 0x23, 0x60, 0x03, 0x00, 0x06, 0x01, 0x5e, 0x20, 0x04,
|
||||
0x40, 0x96, 0x75, 0x20, 0x07, 0x24, 0x24, 0x09, 0x40, 0x20, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x00, 0x2d, 0x00, 0x0d, 0x8d, 0x0a, 0x38, 0x01, 0x00, 0x00,
|
||||
0x31, 0x20, 0x80, 0x0c, 0x68, 0x02, 0x40, 0x2d, 0x80, 0x0d, 0x00, 0x06,
|
||||
0x01, 0x5e, 0x20, 0x04, 0x10, 0x17, 0x17, 0x25, 0x00, 0x00, 0x7e, 0x09,
|
||||
0x31, 0x00, 0x80, 0x0c, 0x68, 0x02, 0xa0, 0x2c, 0x80, 0x04, 0x00, 0x06,
|
||||
0x02, 0x5e, 0x20, 0x04, 0x31, 0x20, 0x80, 0x0c, 0x68, 0x02, 0x60, 0x2c,
|
||||
0x00, 0x0d, 0x00, 0x06, 0x02, 0x5e, 0x20, 0x04, 0x10, 0x20, 0x80, 0x05,
|
||||
0x24, 0x0a, 0x00, 0x20, 0xc4, 0x0f, 0x00, 0x0a, 0x20, 0x01, 0x00, 0x00,
|
||||
0x41, 0x56, 0x5e, 0x20, 0x07, 0x20, 0x1d, 0x09, 0x41, 0x20, 0x80, 0x00,
|
||||
0xe8, 0x3a, 0x40, 0x24, 0x40, 0x0d, 0x8d, 0x3a, 0x28, 0x01, 0x00, 0x00,
|
||||
0x5b, 0xe2, 0x06, 0x20, 0x00, 0x70, 0x83, 0xca, 0x5b, 0x93, 0x01, 0x20,
|
||||
0x00, 0xc8, 0x88, 0xc6, 0x33, 0x00, 0x80, 0x0c, 0x70, 0xe0, 0x06, 0x00,
|
||||
0x82, 0x0e, 0x00, 0x00, 0x00, 0x5e, 0x02, 0x04, 0x33, 0x20, 0x80, 0x0c,
|
||||
0x70, 0x90, 0x01, 0x00, 0x62, 0x02, 0x00, 0x00, 0x00, 0x5e, 0x02, 0x04,
|
||||
0x20, 0x00, 0x01, 0x00, 0x04, 0x00, 0x00, 0x34, 0x00, 0x14, 0x00, 0x0e,
|
||||
0x90, 0xfe, 0xff, 0xff, 0x25, 0x00, 0xa0, 0x00, 0x00, 0x00, 0x00, 0x20,
|
||||
0x00, 0x00, 0x00, 0x0e, 0x10, 0x00, 0x00, 0x00, 0x01, 0x4d, 0x00, 0x20,
|
||||
0x07, 0x7f, 0x05, 0x00, 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,
|
||||
0xc3, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0x04, 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, 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,
|
||||
0x80, 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,
|
||||
0xc0, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x14, 0x00, 0x00, 0x00,
|
||||
0xc0, 0x00, 0x00, 0x00, 0x03, 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, 0x2b, 0x00, 0x00, 0x00, 0x02, 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, 0x2a, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x50, 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,
|
||||
0x2a, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x54, 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, 0x2a, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
|
||||
0x58, 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, 0x60, 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, 0x64, 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,
|
||||
0x68, 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,
|
||||
0x1e, 0x00, 0x00, 0x00, 0x24, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
|
||||
0x80, 0x00, 0x00, 0x00, 0x30, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00,
|
||||
0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
|
||||
0x03, 0x00, 0x00, 0x00, 0x38, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
|
||||
0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00,
|
||||
0x01, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x3c, 0x00, 0x00, 0x00,
|
||||
0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff,
|
||||
0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
|
||||
0x28, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00,
|
||||
0x40, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00,
|
||||
0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
|
||||
0x06, 0x00, 0x00, 0x00, 0x44, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
|
||||
0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00,
|
||||
0x01, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, 0x48, 0x00, 0x00, 0x00,
|
||||
0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff,
|
||||
0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
|
||||
0x28, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x4c, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00,
|
||||
0x19, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x80, 0x00, 0x00, 0x00,
|
||||
0x16, 0x00, 0x00, 0x00, 0x58, 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, 0x00, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x17, 0x00, 0x00, 0x00, 0x8c, 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, 0x07, 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,
|
||||
0x1b, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x50, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x04, 0x00, 0x00, 0x00, 0x10, 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, 0x74, 0x6d, 0x70, 0x00,
|
||||
0x44, 0x41, 0x54, 0x41, 0x5f, 0x54, 0x59, 0x50, 0x45, 0x2a, 0x3b, 0x38,
|
||||
0x00, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
|
||||
0x1a, 0x00, 0x00, 0x00, 0x50, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
|
||||
0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
|
||||
0x10, 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, 0x41, 0x00, 0x00, 0x00, 0x44, 0x41, 0x54, 0x41,
|
||||
0x5f, 0x54, 0x59, 0x50, 0x45, 0x2a, 0x3b, 0x38, 0x00, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
|
||||
0x50, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x10, 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,
|
||||
0x42, 0x00, 0x00, 0x00, 0x44, 0x41, 0x54, 0x41, 0x5f, 0x54, 0x59, 0x50,
|
||||
0x45, 0x2a, 0x3b, 0x38, 0x00, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45,
|
||||
0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x48, 0x00, 0x00, 0x00,
|
||||
0x03, 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, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74, 0x65, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x6e, 0x69, 0x00, 0x00,
|
||||
0x69, 0x6e, 0x74, 0x3b, 0x34, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45,
|
||||
0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x48, 0x00, 0x00, 0x00,
|
||||
0x04, 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, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74, 0x65, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x6e, 0x6a, 0x00, 0x00,
|
||||
0x69, 0x6e, 0x74, 0x3b, 0x34, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45,
|
||||
0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x48, 0x00, 0x00, 0x00,
|
||||
0x05, 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, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74, 0x65, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x6e, 0x6b, 0x00, 0x00,
|
||||
0x69, 0x6e, 0x74, 0x3b, 0x34, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45,
|
||||
0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x48, 0x00, 0x00, 0x00,
|
||||
0x06, 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, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74, 0x65, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x6e, 0x6c, 0x00, 0x00,
|
||||
0x69, 0x6e, 0x74, 0x3b, 0x34, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45,
|
||||
0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x50, 0x00, 0x00, 0x00,
|
||||
0x07, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x5f, 0x5f, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74, 0x65, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x61, 0x6c, 0x70, 0x68,
|
||||
0x61, 0x00, 0x00, 0x00, 0x44, 0x41, 0x54, 0x41, 0x5f, 0x54, 0x59, 0x50,
|
||||
0x45, 0x3b, 0x34, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
|
||||
0x1a, 0x00, 0x00, 0x00, 0x50, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x5f, 0x5f, 0x70, 0x72,
|
||||
0x69, 0x76, 0x61, 0x74, 0x65, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45,
|
||||
0x00, 0x00, 0x00, 0x00, 0x62, 0x65, 0x74, 0x61, 0x00, 0x00, 0x00, 0x00,
|
||||
0x44, 0x41, 0x54, 0x41, 0x5f, 0x54, 0x59, 0x50, 0x45, 0x3b, 0x34, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x0e, 0xe7, 0x4b, 0xdd,
|
||||
0x49, 0x73, 0x32, 0xcf, 0xd7, 0x42, 0x73, 0xa9, 0x0c, 0x00, 0x00, 0x00,
|
||||
0x9c, 0x07, 0x00, 0x00, 0x80, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0x20, 0x00, 0x00, 0x00, 0xcc, 0x00, 0x00, 0x00, 0xd0, 0x03, 0x00, 0x00,
|
||||
0x6d, 0x6d, 0x32, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x32, 0x00,
|
||||
0x01, 0x00, 0x60, 0x00, 0x0c, 0x02, 0xa0, 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, 0x80, 0x2d, 0x20,
|
||||
0x00, 0x7e, 0x0a, 0x05, 0x41, 0x00, 0x00, 0x00, 0x2c, 0x0a, 0xc0, 0x20,
|
||||
0x44, 0x01, 0x00, 0x0a, 0xb8, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
|
||||
0x4c, 0x16, 0xc4, 0x20, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x01, 0x00,
|
||||
0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x80, 0x2f, 0xc0, 0x0f, 0x00, 0x12,
|
||||
0x40, 0x00, 0xb1, 0x00, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0xe0, 0x21,
|
||||
0xc0, 0x00, 0x00, 0x12, 0x80, 0x00, 0xb1, 0x00, 0x40, 0x00, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x60, 0x21, 0xc0, 0x0f, 0x00, 0x12, 0x20, 0x00, 0xb1, 0x00,
|
||||
0x40, 0x00, 0x80, 0x00, 0x28, 0x0a, 0x00, 0x2f, 0xc0, 0x00, 0x00, 0x12,
|
||||
0x60, 0x00, 0xb1, 0x00, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x40, 0x2f,
|
||||
0x80, 0x0f, 0x8d, 0x0a, 0xe0, 0x00, 0x00, 0x00, 0x40, 0x20, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x20, 0x22, 0xe0, 0x01, 0x8d, 0x0a, 0xe4, 0x00, 0x00, 0x00,
|
||||
0x40, 0x96, 0x01, 0x20, 0x07, 0x0d, 0x0b, 0x07, 0x40, 0x96, 0x2d, 0x20,
|
||||
0x07, 0x76, 0x78, 0x07, 0x10, 0x20, 0x80, 0x05, 0x20, 0x0a, 0x00, 0x20,
|
||||
0x40, 0x0f, 0x8d, 0x0a, 0x24, 0x01, 0x00, 0x00, 0x10, 0x20, 0x80, 0x05,
|
||||
0x22, 0x0a, 0x00, 0x20, 0x20, 0x02, 0x8d, 0x0a, 0x18, 0x01, 0x00, 0x00,
|
||||
0x10, 0x00, 0x80, 0x05, 0x22, 0x0a, 0x00, 0x20, 0xa0, 0x01, 0x8d, 0x0a,
|
||||
0x24, 0x01, 0x00, 0x00, 0x02, 0x20, 0x81, 0x00, 0x48, 0x12, 0x20, 0x2d,
|
||||
0xc4, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x00, 0x02, 0x20, 0x81, 0x00,
|
||||
0x4a, 0x12, 0xa0, 0x23, 0xc4, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x00,
|
||||
0x10, 0x00, 0x81, 0x05, 0x22, 0x0a, 0x00, 0x20, 0xc0, 0x0e, 0x8d, 0x0a,
|
||||
0x18, 0x01, 0x00, 0x00, 0x05, 0x20, 0x80, 0x02, 0x42, 0x12, 0x00, 0x20,
|
||||
0xa0, 0x03, 0xb1, 0x12, 0x20, 0x0d, 0xb1, 0x00, 0x22, 0x00, 0xa1, 0x00,
|
||||
0x02, 0x0e, 0x00, 0x20, 0x80, 0x02, 0x00, 0x00, 0x80, 0x02, 0x00, 0x00,
|
||||
0x41, 0x96, 0x2d, 0x20, 0x07, 0x74, 0x76, 0x09, 0x41, 0x20, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x60, 0x22, 0x20, 0x02, 0x8d, 0x0a, 0x24, 0x01, 0x00, 0x00,
|
||||
0x10, 0x00, 0x80, 0x03, 0x24, 0x0a, 0x00, 0x20, 0x1c, 0x01, 0x00, 0x1e,
|
||||
0x00, 0x00, 0x00, 0x00, 0x10, 0x20, 0x80, 0x03, 0x24, 0x0a, 0x00, 0x20,
|
||||
0x1c, 0x01, 0x00, 0x1e, 0x00, 0x00, 0x00, 0x00, 0x40, 0x96, 0x01, 0x20,
|
||||
0xe7, 0x74, 0x74, 0x0d, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x22,
|
||||
0x60, 0x02, 0x8d, 0x0a, 0x40, 0x0f, 0x8d, 0x00, 0x09, 0x00, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x80, 0x2e, 0x80, 0x0e, 0x8d, 0x1e, 0x02, 0x00, 0x02, 0x00,
|
||||
0x09, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x22, 0x60, 0x02, 0x8d, 0x1e,
|
||||
0x02, 0x00, 0x02, 0x00, 0x40, 0x96, 0x75, 0x20, 0x07, 0x74, 0x74, 0x09,
|
||||
0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x22, 0x60, 0x02, 0x8d, 0x0a,
|
||||
0x38, 0x01, 0x00, 0x00, 0x31, 0x00, 0x80, 0x0c, 0x68, 0x02, 0x40, 0x2e,
|
||||
0x80, 0x0e, 0x00, 0x06, 0x02, 0x5e, 0x20, 0x04, 0x31, 0x20, 0x80, 0x0c,
|
||||
0x68, 0x02, 0xa0, 0x22, 0x60, 0x02, 0x00, 0x06, 0x02, 0x5e, 0x20, 0x04,
|
||||
0x41, 0x56, 0x66, 0x20, 0x07, 0x70, 0x72, 0x09, 0x41, 0x20, 0x80, 0x00,
|
||||
0xe8, 0x3a, 0xe0, 0x22, 0xa0, 0x02, 0x8d, 0x3a, 0x2c, 0x01, 0x00, 0x00,
|
||||
0x33, 0x00, 0x80, 0x0c, 0x70, 0x00, 0x07, 0x00, 0x82, 0x0e, 0x00, 0x00,
|
||||
0x02, 0x5e, 0x02, 0x04, 0x33, 0x20, 0x80, 0x0c, 0x70, 0x70, 0x01, 0x00,
|
||||
0x62, 0x02, 0x00, 0x00, 0x02, 0x5e, 0x02, 0x04, 0x20, 0x00, 0x11, 0x00,
|
||||
0x04, 0x00, 0x00, 0x34, 0x00, 0x14, 0x00, 0x0e, 0x80, 0x01, 0x00, 0x00,
|
||||
0x41, 0x96, 0x79, 0x20, 0x07, 0x6e, 0x76, 0x08, 0x41, 0x20, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x20, 0x23, 0x20, 0x02, 0x8d, 0x0a, 0x1c, 0x01, 0x00, 0x00,
|
||||
0x01, 0x00, 0x00, 0x00, 0x2c, 0x1e, 0xc4, 0x2f, 0x00, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x41, 0x80, 0x49, 0x20, 0x00, 0x7e, 0x7e, 0x09,
|
||||
0x40, 0x96, 0x2d, 0x20, 0x07, 0x1b, 0x6e, 0x7e, 0x40, 0x20, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x80, 0x2d, 0x20, 0x03, 0x8d, 0x0a, 0xc4, 0x0f, 0x00, 0x00,
|
||||
0x40, 0x00, 0x00, 0x00, 0x2c, 0x0a, 0xc4, 0x2f, 0xc4, 0x0f, 0x00, 0x1e,
|
||||
0x01, 0x00, 0x01, 0x00, 0x40, 0x96, 0x19, 0x20, 0xe0, 0x22, 0x7e, 0x0d,
|
||||
0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x40, 0x2d, 0xc8, 0x0f, 0x00, 0x0a,
|
||||
0x40, 0x0f, 0x8d, 0x00, 0x09, 0x00, 0x80, 0x00, 0x28, 0x0a, 0x60, 0x23,
|
||||
0x60, 0x03, 0x8d, 0x1e, 0x02, 0x00, 0x02, 0x00, 0x09, 0x20, 0x80, 0x00,
|
||||
0x28, 0x0a, 0x80, 0x2d, 0x80, 0x0d, 0x8d, 0x1e, 0x02, 0x00, 0x02, 0x00,
|
||||
0x09, 0x00, 0x80, 0x00, 0x28, 0x0a, 0x40, 0x24, 0x40, 0x04, 0x8d, 0x1e,
|
||||
0x02, 0x00, 0x02, 0x00, 0x09, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x40, 0x2d,
|
||||
0x40, 0x0d, 0x8d, 0x1e, 0x02, 0x00, 0x02, 0x00, 0x40, 0x96, 0x6d, 0x20,
|
||||
0x07, 0x1b, 0x1b, 0x09, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x80, 0x2d,
|
||||
0x80, 0x0d, 0x8d, 0x0a, 0x30, 0x01, 0x00, 0x00, 0x40, 0x96, 0x71, 0x20,
|
||||
0x07, 0x22, 0x22, 0x09, 0x40, 0x20, 0x80, 0x00, 0x28, 0x0a, 0x40, 0x2d,
|
||||
0x40, 0x0d, 0x8d, 0x0a, 0x34, 0x01, 0x00, 0x00, 0x31, 0x00, 0x80, 0x0c,
|
||||
0x68, 0x02, 0xc0, 0x23, 0x60, 0x03, 0x00, 0x06, 0x00, 0x5e, 0x20, 0x04,
|
||||
0x10, 0x00, 0x80, 0x05, 0x24, 0x0a, 0x00, 0x20, 0xc4, 0x0f, 0x00, 0x0a,
|
||||
0x1c, 0x01, 0x00, 0x00, 0x31, 0x00, 0x80, 0x0c, 0x68, 0x02, 0xe0, 0x2c,
|
||||
0x40, 0x04, 0x00, 0x06, 0x01, 0x5e, 0x20, 0x04, 0x31, 0x20, 0x80, 0x0c,
|
||||
0x68, 0x02, 0x00, 0x24, 0x80, 0x0d, 0x00, 0x06, 0x00, 0x5e, 0x20, 0x04,
|
||||
0x31, 0x20, 0x80, 0x0c, 0x68, 0x02, 0xa0, 0x2c, 0x40, 0x0d, 0x00, 0x06,
|
||||
0x01, 0x5e, 0x20, 0x04, 0x10, 0x20, 0x80, 0x05, 0x24, 0x0a, 0x00, 0x20,
|
||||
0xc4, 0x0f, 0x00, 0x0a, 0x1c, 0x01, 0x00, 0x00, 0x5b, 0x02, 0x07, 0x20,
|
||||
0x00, 0x80, 0x7b, 0xce, 0x5b, 0x73, 0x01, 0x20, 0x00, 0xb8, 0x80, 0xca,
|
||||
0x33, 0x00, 0x80, 0x0c, 0x70, 0x00, 0x07, 0x00, 0x82, 0x0e, 0x00, 0x00,
|
||||
0x02, 0x5e, 0x02, 0x04, 0x33, 0x20, 0x80, 0x0c, 0x70, 0x70, 0x01, 0x00,
|
||||
0x62, 0x02, 0x00, 0x00, 0x02, 0x5e, 0x02, 0x04, 0x20, 0x00, 0x01, 0x00,
|
||||
0x04, 0x00, 0x00, 0x34, 0x00, 0x14, 0x00, 0x0e, 0xa8, 0xfe, 0xff, 0xff,
|
||||
0x25, 0x00, 0xa0, 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00, 0x0e,
|
||||
0x10, 0x00, 0x00, 0x00, 0x01, 0x4d, 0x00, 0x20, 0x07, 0x7f, 0x05, 0x00,
|
||||
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,
|
||||
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0xc3, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x04, 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, 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, 0x80, 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, 0xc0, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x14, 0x00, 0x00, 0x00, 0xc0, 0x00, 0x00, 0x00, 0x03, 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, 0x2b, 0x00, 0x00, 0x00,
|
||||
0x02, 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,
|
||||
0x2a, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x50, 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, 0x2a, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
|
||||
0x54, 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, 0x2a, 0x00, 0x00, 0x00,
|
||||
0x02, 0x00, 0x00, 0x00, 0x58, 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, 0x60, 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,
|
||||
0x64, 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, 0x68, 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, 0x1e, 0x00, 0x00, 0x00, 0x24, 0x00, 0x00, 0x00,
|
||||
0x02, 0x00, 0x00, 0x00, 0x80, 0x00, 0x00, 0x00, 0x30, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
|
||||
0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00,
|
||||
0x01, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x38, 0x00, 0x00, 0x00,
|
||||
0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff,
|
||||
0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
|
||||
0x28, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
|
||||
0x3c, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00,
|
||||
0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
|
||||
0x05, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
|
||||
0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00,
|
||||
0x01, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x44, 0x00, 0x00, 0x00,
|
||||
0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff,
|
||||
0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
|
||||
0x28, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00,
|
||||
0x48, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00,
|
||||
0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0x4c, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
|
||||
0x00, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
|
||||
0x80, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x58, 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,
|
||||
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
|
||||
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x17, 0x00, 0x00, 0x00,
|
||||
0x8c, 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, 0x07, 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, 0x1b, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00,
|
||||
0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
|
||||
0x50, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x10, 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,
|
||||
0x74, 0x6d, 0x70, 0x00, 0x44, 0x41, 0x54, 0x41, 0x5f, 0x54, 0x59, 0x50,
|
||||
0x45, 0x2a, 0x3b, 0x38, 0x00, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45,
|
||||
0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x50, 0x00, 0x00, 0x00,
|
||||
0x01, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x04, 0x00, 0x00, 0x00, 0x10, 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, 0x43, 0x00, 0x00, 0x00,
|
||||
0x44, 0x41, 0x54, 0x41, 0x5f, 0x54, 0x59, 0x50, 0x45, 0x2a, 0x3b, 0x38,
|
||||
0x00, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
|
||||
0x1a, 0x00, 0x00, 0x00, 0x50, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
|
||||
0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
|
||||
0x10, 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, 0x44, 0x00, 0x00, 0x00, 0x44, 0x41, 0x54, 0x41,
|
||||
0x5f, 0x54, 0x59, 0x50, 0x45, 0x2a, 0x3b, 0x38, 0x00, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
|
||||
0x48, 0x00, 0x00, 0x00, 0x03, 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, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74,
|
||||
0x65, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
|
||||
0x6e, 0x69, 0x00, 0x00, 0x69, 0x6e, 0x74, 0x3b, 0x34, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
|
||||
0x48, 0x00, 0x00, 0x00, 0x04, 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, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74,
|
||||
0x65, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
|
||||
0x6e, 0x6a, 0x00, 0x00, 0x69, 0x6e, 0x74, 0x3b, 0x34, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
|
||||
0x48, 0x00, 0x00, 0x00, 0x05, 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, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74,
|
||||
0x65, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
|
||||
0x6e, 0x6b, 0x00, 0x00, 0x69, 0x6e, 0x74, 0x3b, 0x34, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
|
||||
0x48, 0x00, 0x00, 0x00, 0x06, 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, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74,
|
||||
0x65, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
|
||||
0x6e, 0x6c, 0x00, 0x00, 0x69, 0x6e, 0x74, 0x3b, 0x34, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
|
||||
0x50, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0x5f, 0x5f, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74,
|
||||
0x65, 0x00, 0x00, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00,
|
||||
0x61, 0x6c, 0x70, 0x68, 0x61, 0x00, 0x00, 0x00, 0x44, 0x41, 0x54, 0x41,
|
||||
0x5f, 0x54, 0x59, 0x50, 0x45, 0x3b, 0x34, 0x00, 0x4e, 0x4f, 0x4e, 0x45,
|
||||
0x00, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x50, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x08, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
|
||||
0x5f, 0x5f, 0x70, 0x72, 0x69, 0x76, 0x61, 0x74, 0x65, 0x00, 0x00, 0x00,
|
||||
0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00, 0x62, 0x65, 0x74, 0x61,
|
||||
0x00, 0x00, 0x00, 0x00, 0x44, 0x41, 0x54, 0x41, 0x5f, 0x54, 0x59, 0x50,
|
||||
0x45, 0x3b, 0x34, 0x00, 0x4e, 0x4f, 0x4e, 0x45, 0x00, 0x00, 0x00, 0x00
|
||||
};
|
||||
unsigned int __2mm_Gen9core_gen_len = 6804;
|
||||
1642
repos/hello_gpgpu/src/hello_gpgpu/CL/cl.cc
Normal file
1642
repos/hello_gpgpu/src/hello_gpgpu/CL/cl.cc
Normal file
File diff suppressed because it is too large
Load Diff
1934
repos/hello_gpgpu/src/hello_gpgpu/CL/cl.h
Normal file
1934
repos/hello_gpgpu/src/hello_gpgpu/CL/cl.h
Normal file
File diff suppressed because it is too large
Load Diff
1429
repos/hello_gpgpu/src/hello_gpgpu/CL/cl_platform.h
Normal file
1429
repos/hello_gpgpu/src/hello_gpgpu/CL/cl_platform.h
Normal file
File diff suppressed because it is too large
Load Diff
81
repos/hello_gpgpu/src/hello_gpgpu/CL/cl_version.h
Normal file
81
repos/hello_gpgpu/src/hello_gpgpu/CL/cl_version.h
Normal file
@@ -0,0 +1,81 @@
|
||||
/*******************************************************************************
|
||||
* Copyright (c) 2018-2020 The Khronos Group Inc.
|
||||
*
|
||||
* 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.
|
||||
******************************************************************************/
|
||||
|
||||
#ifndef __CL_VERSION_H
|
||||
#define __CL_VERSION_H
|
||||
|
||||
/* Detect which version to target */
|
||||
#if !defined(CL_TARGET_OPENCL_VERSION)
|
||||
#pragma message("cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 300 (OpenCL 3.0)")
|
||||
#define CL_TARGET_OPENCL_VERSION 300
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION != 100 && \
|
||||
CL_TARGET_OPENCL_VERSION != 110 && \
|
||||
CL_TARGET_OPENCL_VERSION != 120 && \
|
||||
CL_TARGET_OPENCL_VERSION != 200 && \
|
||||
CL_TARGET_OPENCL_VERSION != 210 && \
|
||||
CL_TARGET_OPENCL_VERSION != 220 && \
|
||||
CL_TARGET_OPENCL_VERSION != 300
|
||||
#pragma message("cl_version: CL_TARGET_OPENCL_VERSION is not a valid value (100, 110, 120, 200, 210, 220, 300). Defaulting to 300 (OpenCL 3.0)")
|
||||
#undef CL_TARGET_OPENCL_VERSION
|
||||
#define CL_TARGET_OPENCL_VERSION 300
|
||||
#endif
|
||||
|
||||
|
||||
/* OpenCL Version */
|
||||
#if CL_TARGET_OPENCL_VERSION >= 300 && !defined(CL_VERSION_3_0)
|
||||
#define CL_VERSION_3_0 1
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION >= 220 && !defined(CL_VERSION_2_2)
|
||||
#define CL_VERSION_2_2 1
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION >= 210 && !defined(CL_VERSION_2_1)
|
||||
#define CL_VERSION_2_1 1
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION >= 200 && !defined(CL_VERSION_2_0)
|
||||
#define CL_VERSION_2_0 1
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION >= 120 && !defined(CL_VERSION_1_2)
|
||||
#define CL_VERSION_1_2 1
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION >= 110 && !defined(CL_VERSION_1_1)
|
||||
#define CL_VERSION_1_1 1
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION >= 100 && !defined(CL_VERSION_1_0)
|
||||
#define CL_VERSION_1_0 1
|
||||
#endif
|
||||
|
||||
/* Allow deprecated APIs for older OpenCL versions. */
|
||||
#if CL_TARGET_OPENCL_VERSION <= 220 && !defined(CL_USE_DEPRECATED_OPENCL_2_2_APIS)
|
||||
#define CL_USE_DEPRECATED_OPENCL_2_2_APIS
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION <= 210 && !defined(CL_USE_DEPRECATED_OPENCL_2_1_APIS)
|
||||
#define CL_USE_DEPRECATED_OPENCL_2_1_APIS
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION <= 200 && !defined(CL_USE_DEPRECATED_OPENCL_2_0_APIS)
|
||||
#define CL_USE_DEPRECATED_OPENCL_2_0_APIS
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION <= 120 && !defined(CL_USE_DEPRECATED_OPENCL_1_2_APIS)
|
||||
#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION <= 110 && !defined(CL_USE_DEPRECATED_OPENCL_1_1_APIS)
|
||||
#define CL_USE_DEPRECATED_OPENCL_1_1_APIS
|
||||
#endif
|
||||
#if CL_TARGET_OPENCL_VERSION <= 100 && !defined(CL_USE_DEPRECATED_OPENCL_1_0_APIS)
|
||||
#define CL_USE_DEPRECATED_OPENCL_1_0_APIS
|
||||
#endif
|
||||
|
||||
#endif /* __CL_VERSION_H */
|
||||
44
repos/hello_gpgpu/src/hello_gpgpu/main.cc
Normal file
44
repos/hello_gpgpu/src/hello_gpgpu/main.cc
Normal file
@@ -0,0 +1,44 @@
|
||||
#include <base/component.h>
|
||||
#include <base/log.h>
|
||||
#include <base/heap.h>
|
||||
#include <base/allocator_avl.h>
|
||||
#include <dataspace/client.h>
|
||||
#include <hello_gpgpu_session/connection.h>
|
||||
#include "test.h"
|
||||
#include <libc/component.h>
|
||||
#include "CL/cl.h"
|
||||
|
||||
extern int main(int argc, char *argv[]);
|
||||
|
||||
void testvm_construct(Genode::Env &env)
|
||||
{
|
||||
gpgpu::Connection gpgpu(env);
|
||||
|
||||
// allocator
|
||||
Genode::Heap heap(env.ram(), env.rm());
|
||||
Genode::Allocator_avl alloc(&heap);
|
||||
const unsigned int size = 0x10000 * 0x1000;
|
||||
Genode::Ram_dataspace_capability ram_cap = env.ram().alloc(size);
|
||||
Genode::addr_t mapped_base = env.rm().attach(ram_cap);
|
||||
//Genode::addr_t base = Genode::Dataspace_client(ram_cap).phys_addr();
|
||||
alloc.add_range(mapped_base, size);
|
||||
|
||||
// test RPC
|
||||
gpgpu.say_hello();
|
||||
|
||||
// run the test and hope the best
|
||||
//run_gpgpu_test(alloc);
|
||||
|
||||
// run 2mm
|
||||
Libc::with_libc([&] {
|
||||
clInitGenode(alloc);
|
||||
main(0, 0);
|
||||
});
|
||||
|
||||
Genode::log("hello gpgpu completed");
|
||||
}
|
||||
|
||||
void Libc::Component::construct(Libc::Env &env)
|
||||
{
|
||||
testvm_construct(env);
|
||||
}
|
||||
402
repos/hello_gpgpu/src/hello_gpgpu/polybench.cc
Normal file
402
repos/hello_gpgpu/src/hello_gpgpu/polybench.cc
Normal file
@@ -0,0 +1,402 @@
|
||||
/**
|
||||
* polybench.c: This file is part of the PolyBench/C 3.2 test suite.
|
||||
*
|
||||
*
|
||||
* Contact: Louis-Noel Pouchet <pouchet@cse.ohio-state.edu>
|
||||
* Web address: http://polybench.sourceforge.net
|
||||
*/
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
#include <stdlib.h>
|
||||
#include <unistd.h>
|
||||
#include <assert.h>
|
||||
#include <time.h>
|
||||
#include <sys/time.h>
|
||||
#include <sys/resource.h>
|
||||
#include <sched.h>
|
||||
#include <math.h>
|
||||
#ifdef _OPENMP
|
||||
# include <omp.h>
|
||||
#endif
|
||||
|
||||
/* By default, collect PAPI counters on thread 0. */
|
||||
#ifndef POLYBENCH_THREAD_MONITOR
|
||||
# define POLYBENCH_THREAD_MONITOR 0
|
||||
#endif
|
||||
|
||||
/* Total LLC cache size. By default 32+MB.. */
|
||||
#ifndef POLYBENCH_CACHE_SIZE_KB
|
||||
# define POLYBENCH_CACHE_SIZE_KB 32770
|
||||
#endif
|
||||
|
||||
|
||||
int polybench_papi_counters_threadid = POLYBENCH_THREAD_MONITOR;
|
||||
double polybench_program_total_flops = 0;
|
||||
|
||||
#ifdef POLYBENCH_PAPI
|
||||
# include <papi.h>
|
||||
# define POLYBENCH_MAX_NB_PAPI_COUNTERS 96
|
||||
char* _polybench_papi_eventlist[] = {
|
||||
#include "papi_counters.list"
|
||||
NULL
|
||||
};
|
||||
int polybench_papi_eventset;
|
||||
int polybench_papi_eventlist[POLYBENCH_MAX_NB_PAPI_COUNTERS];
|
||||
long_long polybench_papi_values[POLYBENCH_MAX_NB_PAPI_COUNTERS];
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
/* Timer code (gettimeofday). */
|
||||
double polybench_t_start, polybench_t_end;
|
||||
/* Timer code (RDTSC). */
|
||||
unsigned long long int polybench_c_start, polybench_c_end;
|
||||
|
||||
static
|
||||
double rtclock()
|
||||
{
|
||||
#ifdef POLYBENCH_TIME
|
||||
struct timeval Tp;
|
||||
int stat;
|
||||
stat = gettimeofday (&Tp, NULL);
|
||||
if (stat != 0)
|
||||
printf ("Error return from gettimeofday: %d", stat);
|
||||
return ((double)Tp.tv_sec + (double)Tp.tv_usec * 1.0e-6);
|
||||
#else
|
||||
return 0;
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
#ifdef POLYBENCH_CYCLE_ACCURATE_TIMER
|
||||
static
|
||||
unsigned long long int rdtsc()
|
||||
{
|
||||
unsigned long long int ret = 0;
|
||||
unsigned int cycles_lo;
|
||||
unsigned int cycles_hi;
|
||||
__asm__ volatile ("RDTSC" : "=a" (cycles_lo), "=d" (cycles_hi));
|
||||
ret = (unsigned long long int)cycles_hi << 32 | cycles_lo;
|
||||
|
||||
return ret;
|
||||
}
|
||||
#endif
|
||||
|
||||
void polybench_flush_cache()
|
||||
{
|
||||
int cs = POLYBENCH_CACHE_SIZE_KB * 1024 / sizeof(double);
|
||||
double* flush = (double*) calloc (cs, sizeof(double));
|
||||
int i;
|
||||
double tmp = 0.0;
|
||||
#ifdef _OPENMP
|
||||
#pragma omp parallel for
|
||||
#endif
|
||||
for (i = 0; i < cs; i++)
|
||||
tmp += flush[i];
|
||||
assert (tmp <= 10.0);
|
||||
free (flush);
|
||||
}
|
||||
|
||||
|
||||
#ifdef POLYBENCH_LINUX_FIFO_SCHEDULER
|
||||
void polybench_linux_fifo_scheduler()
|
||||
{
|
||||
/* Use FIFO scheduler to limit OS interference. Program must be run
|
||||
as root, and this works only for Linux kernels. */
|
||||
struct sched_param schedParam;
|
||||
schedParam.sched_priority = sched_get_priority_max (SCHED_FIFO);
|
||||
sched_setscheduler (0, SCHED_FIFO, &schedParam);
|
||||
}
|
||||
|
||||
|
||||
void polybench_linux_standard_scheduler()
|
||||
{
|
||||
/* Restore to standard scheduler policy. */
|
||||
struct sched_param schedParam;
|
||||
schedParam.sched_priority = sched_get_priority_max (SCHED_OTHER);
|
||||
sched_setscheduler (0, SCHED_OTHER, &schedParam);
|
||||
}
|
||||
#endif
|
||||
|
||||
#ifdef POLYBENCH_PAPI
|
||||
|
||||
static
|
||||
void test_fail(char *file, int line, char *call, int retval)
|
||||
{
|
||||
char buf[128];
|
||||
|
||||
memset(buf, '\0', sizeof(buf));
|
||||
if (retval != 0)
|
||||
fprintf (stdout,"%-40s FAILED\nLine # %d\n", file, line);
|
||||
else
|
||||
{
|
||||
fprintf (stdout,"%-40s SKIPPED\n", file);
|
||||
fprintf (stdout,"Line # %d\n", line);
|
||||
}
|
||||
if (retval == PAPI_ESYS)
|
||||
{
|
||||
sprintf (buf, "System error in %s", call);
|
||||
perror (buf);
|
||||
}
|
||||
else if (retval > 0)
|
||||
fprintf (stdout,"Error: %s\n", call);
|
||||
else if (retval == 0)
|
||||
fprintf (stdout,"Error: %s\n", call);
|
||||
else
|
||||
{
|
||||
char errstring[PAPI_MAX_STR_LEN];
|
||||
PAPI_perror (retval, errstring, PAPI_MAX_STR_LEN);
|
||||
fprintf (stdout,"Error in %s: %s\n", call, errstring);
|
||||
}
|
||||
fprintf (stdout,"\n");
|
||||
if (PAPI_is_initialized ())
|
||||
PAPI_shutdown ();
|
||||
exit (1);
|
||||
}
|
||||
|
||||
|
||||
void polybench_papi_init()
|
||||
{
|
||||
# ifdef _OPENMP
|
||||
#pragma omp parallel
|
||||
{
|
||||
#pragma omp master
|
||||
{
|
||||
if (omp_get_max_threads () < polybench_papi_counters_threadid)
|
||||
polybench_papi_counters_threadid = omp_get_max_threads () - 1;
|
||||
}
|
||||
#pragma omp barrier
|
||||
|
||||
if (omp_get_thread_num () == polybench_papi_counters_threadid)
|
||||
{
|
||||
# endif
|
||||
int retval;
|
||||
polybench_papi_eventset = PAPI_NULL;
|
||||
if ((retval = PAPI_library_init (PAPI_VER_CURRENT)) != PAPI_VER_CURRENT)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_library_init", retval);
|
||||
if ((retval = PAPI_create_eventset (&polybench_papi_eventset))
|
||||
!= PAPI_OK)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_create_eventset", retval);
|
||||
int k;
|
||||
for (k = 0; _polybench_papi_eventlist[k]; ++k)
|
||||
{
|
||||
if ((retval =
|
||||
PAPI_event_name_to_code (_polybench_papi_eventlist[k],
|
||||
&(polybench_papi_eventlist[k])))
|
||||
!= PAPI_OK)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_event_name_to_code", retval);
|
||||
}
|
||||
polybench_papi_eventlist[k] = 0;
|
||||
|
||||
|
||||
# ifdef _OPENMP
|
||||
}
|
||||
}
|
||||
#pragma omp barrier
|
||||
# endif
|
||||
}
|
||||
|
||||
|
||||
void polybench_papi_close()
|
||||
{
|
||||
# ifdef _OPENMP
|
||||
#pragma omp parallel
|
||||
{
|
||||
if (omp_get_thread_num () == polybench_papi_counters_threadid)
|
||||
{
|
||||
# endif
|
||||
int retval;
|
||||
if ((retval = PAPI_destroy_eventset (&polybench_papi_eventset))
|
||||
!= PAPI_OK)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_destroy_eventset", retval);
|
||||
if (PAPI_is_initialized ())
|
||||
PAPI_shutdown ();
|
||||
# ifdef _OPENMP
|
||||
}
|
||||
}
|
||||
#pragma omp barrier
|
||||
# endif
|
||||
}
|
||||
|
||||
int polybench_papi_start_counter(int evid)
|
||||
{
|
||||
# ifndef POLYBENCH_NO_FLUSH_CACHE
|
||||
polybench_flush_cache();
|
||||
# endif
|
||||
|
||||
# ifdef _OPENMP
|
||||
# pragma omp parallel
|
||||
{
|
||||
if (omp_get_thread_num () == polybench_papi_counters_threadid)
|
||||
{
|
||||
# endif
|
||||
|
||||
int retval = 1;
|
||||
char descr[PAPI_MAX_STR_LEN];
|
||||
PAPI_event_info_t evinfo;
|
||||
PAPI_event_code_to_name (polybench_papi_eventlist[evid], descr);
|
||||
if (PAPI_add_event (polybench_papi_eventset,
|
||||
polybench_papi_eventlist[evid]) != PAPI_OK)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_add_event", 1);
|
||||
if (PAPI_get_event_info (polybench_papi_eventlist[evid], &evinfo)
|
||||
!= PAPI_OK)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_get_event_info", retval);
|
||||
if ((retval = PAPI_start (polybench_papi_eventset)) != PAPI_OK)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_start", retval);
|
||||
# ifdef _OPENMP
|
||||
}
|
||||
}
|
||||
#pragma omp barrier
|
||||
# endif
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
||||
void polybench_papi_stop_counter(int evid)
|
||||
{
|
||||
# ifdef _OPENMP
|
||||
# pragma omp parallel
|
||||
{
|
||||
if (omp_get_thread_num () == polybench_papi_counters_threadid)
|
||||
{
|
||||
# endif
|
||||
int retval;
|
||||
long_long values[1];
|
||||
values[0] = 0;
|
||||
if ((retval = PAPI_read (polybench_papi_eventset, &values[0]))
|
||||
!= PAPI_OK)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_read", retval);
|
||||
|
||||
if ((retval = PAPI_stop (polybench_papi_eventset, NULL)) != PAPI_OK)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_stop", retval);
|
||||
|
||||
polybench_papi_values[evid] = values[0];
|
||||
|
||||
if ((retval = PAPI_remove_event
|
||||
(polybench_papi_eventset,
|
||||
polybench_papi_eventlist[evid])) != PAPI_OK)
|
||||
test_fail (__FILE__, __LINE__, "PAPI_remove_event", retval);
|
||||
# ifdef _OPENMP
|
||||
}
|
||||
}
|
||||
#pragma omp barrier
|
||||
# endif
|
||||
}
|
||||
|
||||
|
||||
void polybench_papi_print()
|
||||
{
|
||||
int verbose = 0;
|
||||
# ifdef _OPENMP
|
||||
# pragma omp parallel
|
||||
{
|
||||
if (omp_get_thread_num() == polybench_papi_counters_threadid)
|
||||
{
|
||||
#ifdef POLYBENCH_PAPI_VERBOSE
|
||||
verbose = 1;
|
||||
#endif
|
||||
if (verbose)
|
||||
printf ("On thread %d:\n", polybench_papi_counters_threadid);
|
||||
#endif
|
||||
int evid;
|
||||
for (evid = 0; polybench_papi_eventlist[evid] != 0; ++evid)
|
||||
{
|
||||
if (verbose)
|
||||
printf ("%s=", _polybench_papi_eventlist[evid]);
|
||||
printf ("%llu ", polybench_papi_values[evid]);
|
||||
if (verbose)
|
||||
printf ("\n");
|
||||
}
|
||||
printf ("\n");
|
||||
# ifdef _OPENMP
|
||||
}
|
||||
}
|
||||
#pragma omp barrier
|
||||
# endif
|
||||
}
|
||||
|
||||
#endif
|
||||
/* ! POLYBENCH_PAPI */
|
||||
|
||||
void polybench_prepare_instruments()
|
||||
{
|
||||
#ifndef POLYBENCH_NO_FLUSH_CACHE
|
||||
polybench_flush_cache ();
|
||||
#endif
|
||||
#ifdef POLYBENCH_LINUX_FIFO_SCHEDULER
|
||||
polybench_linux_fifo_scheduler ();
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
void polybench_timer_start()
|
||||
{
|
||||
polybench_prepare_instruments ();
|
||||
#ifndef POLYBENCH_CYCLE_ACCURATE_TIMER
|
||||
polybench_t_start = rtclock ();
|
||||
#else
|
||||
polybench_c_start = rdtsc ();
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
void polybench_timer_stop()
|
||||
{
|
||||
#ifndef POLYBENCH_CYCLE_ACCURATE_TIMER
|
||||
polybench_t_end = rtclock ();
|
||||
#else
|
||||
polybench_c_end = rdtsc ();
|
||||
#endif
|
||||
#ifdef POLYBENCH_LINUX_FIFO_SCHEDULER
|
||||
polybench_linux_standard_scheduler ();
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
void polybench_timer_print()
|
||||
{
|
||||
#ifdef POLYBENCH_GFLOPS
|
||||
if (__polybench_program_total_flops == 0)
|
||||
{
|
||||
printf ("[PolyBench][WARNING] Program flops not defined, use polybench_set_program_flops(value)\n");
|
||||
printf ("%0.6lf\n", polybench_t_end - polybench_t_start);
|
||||
}
|
||||
else
|
||||
printf ("%0.2lf\n",
|
||||
(__polybench_program_total_flops /
|
||||
(double)(polybench_t_end - polybench_t_start)) / 1000000000);
|
||||
#else
|
||||
# ifndef POLYBENCH_CYCLE_ACCURATE_TIMER
|
||||
printf ("%0.6f\n", polybench_t_end - polybench_t_start);
|
||||
# else
|
||||
printf ("%Ld\n", polybench_c_end - polybench_c_start);
|
||||
# endif
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
|
||||
static
|
||||
void *
|
||||
xmalloc (size_t num)
|
||||
{
|
||||
void* newA = NULL;
|
||||
int ret = posix_memalign (&newA, 32, num);
|
||||
if (! newA || ret)
|
||||
{
|
||||
fprintf (stderr, "[PolyBench] posix_memalign: cannot allocate memory");
|
||||
exit (1);
|
||||
}
|
||||
return newA;
|
||||
}
|
||||
|
||||
|
||||
void* polybench_alloc_data(unsigned long long int n, int elt_size)
|
||||
{
|
||||
/// FIXME: detect overflow!
|
||||
size_t val = n;
|
||||
val *= elt_size;
|
||||
void* ret = xmalloc (val);
|
||||
|
||||
return ret;
|
||||
}
|
||||
202
repos/hello_gpgpu/src/hello_gpgpu/polybench.h
Normal file
202
repos/hello_gpgpu/src/hello_gpgpu/polybench.h
Normal file
@@ -0,0 +1,202 @@
|
||||
/**
|
||||
* polybench.h: This file is part of the PolyBench/C 3.2 test suite.
|
||||
*
|
||||
*
|
||||
* Contact: Louis-Noel Pouchet <pouchet@cse.ohio-state.edu>
|
||||
* Web address: http://polybench.sourceforge.net
|
||||
*/
|
||||
/*
|
||||
* Polybench header for instrumentation.
|
||||
*
|
||||
* Programs must be compiled with `-I utilities utilities/polybench.c'
|
||||
*
|
||||
* Optionally, one can define:
|
||||
*
|
||||
* -DPOLYBENCH_TIME, to report the execution time,
|
||||
* OR (exclusive):
|
||||
* -DPOLYBENCH_PAPI, to use PAPI H/W counters (defined in polybench.c)
|
||||
*
|
||||
*
|
||||
* See README or utilities/polybench.c for additional options.
|
||||
*
|
||||
*/
|
||||
#ifndef POLYBENCH_H
|
||||
# define POLYBENCH_H
|
||||
|
||||
# include <stdlib.h>
|
||||
|
||||
/* Array padding. By default, none is used. */
|
||||
# ifndef POLYBENCH_PADDING_FACTOR
|
||||
/* default: */
|
||||
# define POLYBENCH_PADDING_FACTOR 0
|
||||
# endif
|
||||
|
||||
|
||||
/* C99 arrays in function prototype. By default, do not use. */
|
||||
# ifdef POLYBENCH_USE_C99_PROTO
|
||||
# define POLYBENCH_C99_SELECT(x,y) y
|
||||
# else
|
||||
/* default: */
|
||||
# define POLYBENCH_C99_SELECT(x,y) x
|
||||
# endif
|
||||
|
||||
|
||||
/* Scalar loop bounds in SCoPs. By default, use parametric loop bounds. */
|
||||
# ifdef POLYBENCH_USE_SCALAR_LB
|
||||
# define POLYBENCH_LOOP_BOUND(x,y) x
|
||||
# else
|
||||
/* default: */
|
||||
# define POLYBENCH_LOOP_BOUND(x,y) y
|
||||
# endif
|
||||
|
||||
|
||||
/* Macros to reference an array. Generic for heap and stack arrays
|
||||
(C99). Each array dimensionality has his own macro, to be used at
|
||||
declaration or as a function argument.
|
||||
Example:
|
||||
int b[x] => POLYBENCH_1D_ARRAY(b, x)
|
||||
int A[N][N] => POLYBENCH_2D_ARRAY(A, N, N)
|
||||
*/
|
||||
# ifndef POLYBENCH_STACK_ARRAYS
|
||||
# define POLYBENCH_ARRAY(x) *x
|
||||
# define POLYBENCH_FREE_ARRAY(x) free((void*)x);
|
||||
# define POLYBENCH_DECL_VAR(x) (*x)
|
||||
# else
|
||||
# define POLYBENCH_ARRAY(x) x
|
||||
# define POLYBENCH_FREE_ARRAY(x)
|
||||
# define POLYBENCH_DECL_VAR(x) x
|
||||
# endif
|
||||
/* Macros for using arrays in the function prototypes. */
|
||||
# define POLYBENCH_1D(var, dim1,ddim1) var[POLYBENCH_C99_SELECT(dim1,ddim1) + POLYBENCH_PADDING_FACTOR]
|
||||
# define POLYBENCH_2D(var, dim1, dim2, ddim1, ddim2) var[POLYBENCH_C99_SELECT(dim1,ddim1) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim2,ddim2) + POLYBENCH_PADDING_FACTOR]
|
||||
# define POLYBENCH_3D(var, dim1, dim2, dim3, ddim1, ddim2, ddim3) var[POLYBENCH_C99_SELECT(dim1,ddim1) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim2,ddim2) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim3,ddim3) + POLYBENCH_PADDING_FACTOR]
|
||||
# define POLYBENCH_4D(var, dim1, dim2, dim3, dim4, ddim1, ddim2, ddim3, ddim4) var[POLYBENCH_C99_SELECT(dim1,ddim1) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim2,ddim2) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim3,ddim3) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim4,ddim4) + POLYBENCH_PADDING_FACTOR]
|
||||
# define POLYBENCH_5D(var, dim1, dim2, dim3, dim4, dim5, ddim1, ddim2, ddim3, ddim4, ddim5) var[POLYBENCH_C99_SELECT(dim1,ddim1) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim2,ddim2) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim3,ddim3) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim4,ddim4) + POLYBENCH_PADDING_FACTOR][POLYBENCH_C99_SELECT(dim5,ddim5) + POLYBENCH_PADDING_FACTOR]
|
||||
|
||||
|
||||
/* Macros to allocate heap arrays.
|
||||
Example:
|
||||
polybench_alloc_2d_array(N, M, double) => allocates N x M x sizeof(double)
|
||||
and returns a pointer to the 2d array
|
||||
*/
|
||||
# define POLYBENCH_ALLOC_1D_ARRAY(n1, type) \
|
||||
(type(*)[n1 + POLYBENCH_PADDING_FACTOR])polybench_alloc_data (n1 + POLYBENCH_PADDING_FACTOR, sizeof(type))
|
||||
# define POLYBENCH_ALLOC_2D_ARRAY(n1, n2, type) \
|
||||
(type(*)[n1 + POLYBENCH_PADDING_FACTOR][n2 + POLYBENCH_PADDING_FACTOR])polybench_alloc_data ((n1 + POLYBENCH_PADDING_FACTOR) * (n2 + POLYBENCH_PADDING_FACTOR), sizeof(type))
|
||||
# define POLYBENCH_ALLOC_3D_ARRAY(n1, n2, n3, type) \
|
||||
(type(*)[n1 + POLYBENCH_PADDING_FACTOR][n2 + POLYBENCH_PADDING_FACTOR][n3 + POLYBENCH_PADDING_FACTOR])polybench_alloc_data ((n1 + POLYBENCH_PADDING_FACTOR) * (n2 + POLYBENCH_PADDING_FACTOR) * (n3 + POLYBENCH_PADDING_FACTOR), sizeof(type))
|
||||
# define POLYBENCH_ALLOC_4D_ARRAY(n1, n2, n3, n4, type) \
|
||||
(type(*)[n1 + POLYBENCH_PADDING_FACTOR][n2 + POLYBENCH_PADDING_FACTOR][n3 + POLYBENCH_PADDING_FACTOR][n4 + POLYBENCH_PADDING_FACTOR])polybench_alloc_data ((n1 + POLYBENCH_PADDING_FACTOR) * (n2 + POLYBENCH_PADDING_FACTOR) * (n3 + POLYBENCH_PADDING_FACTOR) * (n4 + POLYBENCH_PADDING_FACTOR), sizeof(type))
|
||||
# define POLYBENCH_ALLOC_5D_ARRAY(n1, n2, n3, n4, n5, type) \
|
||||
(type(*)[n1 + POLYBENCH_PADDING_FACTOR][n2 + POLYBENCH_PADDING_FACTOR][n3 + POLYBENCH_PADDING_FACTOR][n4 + POLYBENCH_PADDING_FACTOR][n5 + POLYBENCH_PADDING_FACTOR])polybench_alloc_data ((n1 + POLYBENCH_PADDING_FACTOR) * (n2 + POLYBENCH_PADDING_FACTOR) * (n3 + POLYBENCH_PADDING_FACTOR) * (n4 + POLYBENCH_PADDING_FACTOR) * (n5 + POLYBENCH_PADDING_FACTOR), sizeof(type))
|
||||
|
||||
/* Macros for array declaration. */
|
||||
# ifndef POLYBENCH_STACK_ARRAYS
|
||||
# define POLYBENCH_1D_ARRAY_DECL(var, type, dim1, ddim1) \
|
||||
type POLYBENCH_1D(POLYBENCH_DECL_VAR(var), dim1, ddim1); \
|
||||
var = POLYBENCH_ALLOC_1D_ARRAY(POLYBENCH_C99_SELECT(dim1, ddim1), type);
|
||||
# define POLYBENCH_2D_ARRAY_DECL(var, type, dim1, dim2, ddim1, ddim2) \
|
||||
type POLYBENCH_2D(POLYBENCH_DECL_VAR(var), dim1, dim2, ddim1, ddim2); \
|
||||
var = POLYBENCH_ALLOC_2D_ARRAY(POLYBENCH_C99_SELECT(dim1, ddim1), POLYBENCH_C99_SELECT(dim2, ddim2), type);
|
||||
# define POLYBENCH_3D_ARRAY_DECL(var, type, dim1, dim2, dim3, ddim1, ddim2, ddim3) \
|
||||
type POLYBENCH_3D(POLYBENCH_DECL_VAR(var), dim1, dim2, dim3, ddim1, ddim2, ddim3); \
|
||||
var = POLYBENCH_ALLOC_3D_ARRAY(POLYBENCH_C99_SELECT(dim1, ddim1), POLYBENCH_C99_SELECT(dim2, ddim2), POLYBENCH_C99_SELECT(dim3, ddim3), type);
|
||||
# define POLYBENCH_4D_ARRAY_DECL(var, type, dim1, dim2, dim3, dim4, ddim1, ddim2, ddim3, ddim4) \
|
||||
type POLYBENCH_4D(POLYBENCH_DECL_VAR(var), dim1, dim2, ,dim3, dim4, ddim1, ddim2, ddim3, ddim4); \
|
||||
var = POLYBENCH_ALLOC_4D_ARRAY(POLYBENCH_C99_SELECT(dim1, ddim1), POLYBENCH_C99_SELECT(dim2, ddim2), POLYBENCH_C99_SELECT(dim3, ddim3), POLYBENCH_C99_SELECT(dim4, ddim4), type);
|
||||
# define POLYBENCH_5D_ARRAY_DECL(var, type, dim1, dim2, dim3, dim4, dim5, ddim1, ddim2, ddim3, ddim4, ddim5) \
|
||||
type POLYBENCH_5D(POLYBENCH_DECL_VAR(var), dim1, dim2, dim3, dim4, dim5, ddim1, ddim2, ddim3, ddim4, ddim5); \
|
||||
var = POLYBENCH_ALLOC_5D_ARRAY(POLYBENCH_C99_SELECT(dim1, ddim1), POLYBENCH_C99_SELECT(dim2, ddim2), POLYBENCH_C99_SELECT(dim3, ddim3), POLYBENCH_C99_SELECT(dim4, ddim4), POLYBENCH_C99_SELECT(dim5, ddim5), type);
|
||||
# else
|
||||
# define POLYBENCH_1D_ARRAY_DECL(var, type, dim1, ddim1) \
|
||||
type POLYBENCH_1D(POLYBENCH_DECL_VAR(var), dim1, ddim1);
|
||||
# define POLYBENCH_2D_ARRAY_DECL(var, type, dim1, dim2, ddim1, ddim2) \
|
||||
type POLYBENCH_2D(POLYBENCH_DECL_VAR(var), dim1, dim2, ddim1, ddim2);
|
||||
# define POLYBENCH_3D_ARRAY_DECL(var, type, dim1, dim2, dim3, ddim1, ddim2, ddim3) \
|
||||
type POLYBENCH_3D(POLYBENCH_DECL_VAR(var), dim1, dim2, dim3, ddim1, ddim2, ddim3);
|
||||
# define POLYBENCH_4D_ARRAY_DECL(var, type, dim1, dim2, dim3, dim4, ddim1, ddim2, ddim3, ddim4) \
|
||||
type POLYBENCH_4D(POLYBENCH_DECL_VAR(var), dim1, dim2, dim3, dim4, ddim1, ddim2, ddim3, ddim4);
|
||||
# define POLYBENCH_5D_ARRAY_DECL(var, type, dim1, dim2, dim3, dim4, dim5, ddim1, ddim2, ddim3, ddim4, ddim5) \
|
||||
type POLYBENCH_5D(POLYBENCH_DECL_VAR(var), dim1, dim2, dim3, dim4, dim5, ddim1, ddim2, ddim3, ddim4, ddim5);
|
||||
# endif
|
||||
|
||||
|
||||
/* Dead-code elimination macros. Use argc/argv for the run-time check. */
|
||||
# ifndef POLYBENCH_DUMP_ARRAYS
|
||||
# define POLYBENCH_DCE_ONLY_CODE if (argc > 42 && ! strcmp(argv[0], ""))
|
||||
# else
|
||||
# define POLYBENCH_DCE_ONLY_CODE
|
||||
# endif
|
||||
|
||||
# define polybench_prevent_dce(func) \
|
||||
POLYBENCH_DCE_ONLY_CODE \
|
||||
func
|
||||
|
||||
|
||||
/* Performance-related instrumentation. See polybench.c */
|
||||
# define polybench_start_instruments
|
||||
# define polybench_stop_instruments
|
||||
# define polybench_print_instruments
|
||||
|
||||
|
||||
/* PAPI support. */
|
||||
# ifdef POLYBENCH_PAPI
|
||||
extern const unsigned int polybench_papi_eventlist[];
|
||||
# undef polybench_start_instruments
|
||||
# undef polybench_stop_instruments
|
||||
# undef polybench_print_instruments
|
||||
# define polybench_set_papi_thread_report(x) \
|
||||
polybench_papi_counters_threadid = x;
|
||||
# define polybench_start_instruments \
|
||||
polybench_prepare_instruments(); \
|
||||
polybench_papi_init(); \
|
||||
int evid; \
|
||||
for (evid = 0; polybench_papi_eventlist[evid] != 0; evid++) \
|
||||
{ \
|
||||
if (polybench_papi_start_counter(evid)) \
|
||||
continue; \
|
||||
|
||||
# define polybench_stop_instruments \
|
||||
polybench_papi_stop_counter(evid); \
|
||||
} \
|
||||
polybench_papi_close(); \
|
||||
|
||||
# define polybench_print_instruments polybench_papi_print();
|
||||
# endif
|
||||
|
||||
|
||||
/* Timing support. */
|
||||
# if defined(POLYBENCH_TIME) || defined(POLYBENCH_GFLOPS)
|
||||
# undef polybench_start_instruments
|
||||
# undef polybench_stop_instruments
|
||||
# undef polybench_print_instruments
|
||||
# define polybench_start_instruments polybench_timer_start();
|
||||
# define polybench_stop_instruments polybench_timer_stop();
|
||||
# define polybench_print_instruments polybench_timer_print();
|
||||
extern double polybench_program_total_flops;
|
||||
extern void polybench_timer_start();
|
||||
extern void polybench_timer_stop();
|
||||
extern void polybench_timer_print();
|
||||
# endif
|
||||
|
||||
/* Function declaration. */
|
||||
# ifdef POLYBENCH_TIME
|
||||
extern void polybench_timer_start();
|
||||
extern void polybench_timer_stop();
|
||||
extern void polybench_timer_print();
|
||||
# endif
|
||||
|
||||
# ifdef POLYBENCH_PAPI
|
||||
extern void polybench_prepare_instruments();
|
||||
extern int polybench_papi_start_counter(int evid);
|
||||
extern void polybench_papi_stop_counter(int evid);
|
||||
extern void polybench_papi_init();
|
||||
extern void polybench_papi_close();
|
||||
extern void polybench_papi_print();
|
||||
# endif
|
||||
|
||||
/* Function prototypes. */
|
||||
extern void* polybench_alloc_data(unsigned long long int n, int elt_size);
|
||||
|
||||
|
||||
#endif /* !POLYBENCH_H */
|
||||
36
repos/hello_gpgpu/src/hello_gpgpu/polybenchUtilFuncts.h
Normal file
36
repos/hello_gpgpu/src/hello_gpgpu/polybenchUtilFuncts.h
Normal file
@@ -0,0 +1,36 @@
|
||||
//polybenchUtilFuncts.h
|
||||
//Scott Grauer-Gray (sgrauerg@gmail.com)
|
||||
//Functions used across codes
|
||||
|
||||
#ifndef POLYBENCH_UTIL_FUNCTS_H
|
||||
#define POLYBENCH_UTIL_FUNCTS_H
|
||||
|
||||
//define a small float value
|
||||
#define SMALL_FLOAT_VAL 0.00000001f
|
||||
|
||||
double absVal(double a)
|
||||
{
|
||||
if(a < 0)
|
||||
{
|
||||
return (a * -1);
|
||||
}
|
||||
else
|
||||
{
|
||||
return a;
|
||||
}
|
||||
}
|
||||
|
||||
double percentDiff(double val1, double val2)
|
||||
{
|
||||
if ((absVal(val1) < 0.01) && (absVal(val2) < 0.01))
|
||||
{
|
||||
return 0.0f;
|
||||
}
|
||||
|
||||
else
|
||||
{
|
||||
return 100.0f * (absVal(absVal(val1 - val2) / absVal(val1 + SMALL_FLOAT_VAL)));
|
||||
}
|
||||
}
|
||||
|
||||
#endif //POLYBENCH_UTIL_FUNCTS_H
|
||||
5
repos/hello_gpgpu/src/hello_gpgpu/target.mk
Normal file
5
repos/hello_gpgpu/src/hello_gpgpu/target.mk
Normal file
@@ -0,0 +1,5 @@
|
||||
TARGET = hello_gpgpu
|
||||
SRC_CC = main.cc test.cc CL/cl.cc 2mm.cc
|
||||
LIBS = base libc
|
||||
|
||||
CC_CXX_WARN_STRICT =
|
||||
251
repos/hello_gpgpu/src/hello_gpgpu/test.cc
Normal file
251
repos/hello_gpgpu/src/hello_gpgpu/test.cc
Normal file
@@ -0,0 +1,251 @@
|
||||
#include <base/log.h>
|
||||
#include <base/allocator_avl.h>
|
||||
#define CL_TARGET_OPENCL_VERSION 100
|
||||
#include "CL/cl.h"
|
||||
|
||||
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
|
||||
};
|
||||
static unsigned int test_Gen9core_gen_len = 1568;
|
||||
|
||||
#define ELEMENTS 4096
|
||||
void run_gpgpu_test(Genode::Allocator_avl& alloc)
|
||||
{
|
||||
clInitGenode(alloc);
|
||||
const int num = 0x42;
|
||||
uint32_t* m_in;
|
||||
volatile uint32_t* m_out;
|
||||
|
||||
// allocate buffers
|
||||
m_in = (uint32_t*)alloc.alloc(ELEMENTS * sizeof(uint32_t));
|
||||
m_out = (volatile uint32_t*)alloc.alloc(ELEMENTS * sizeof(uint32_t));
|
||||
|
||||
for(int i = 0; i < ELEMENTS; i++)
|
||||
{
|
||||
m_in[i] = num;
|
||||
m_out[i] = 0;
|
||||
}
|
||||
|
||||
cl_platform_id platform_id;
|
||||
cl_device_id device_id;
|
||||
cl_uint num_devices;
|
||||
cl_uint num_platforms;
|
||||
cl_int errcode;
|
||||
cl_context clContext;
|
||||
cl_kernel clKernel;
|
||||
cl_command_queue clCommandQue;
|
||||
cl_program clProgram;
|
||||
cl_mem clInBuff;
|
||||
cl_mem clOutBuff;
|
||||
|
||||
// init opencl stuff
|
||||
errcode = clGetPlatformIDs(1, &platform_id, &num_platforms);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in number of platforms");
|
||||
errcode = clGetDeviceIDs( platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, &num_devices);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in number of devices");
|
||||
clContext = clCreateContext( NULL, 1, &device_id, NULL, NULL, &errcode);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in creating context");
|
||||
clCommandQue = clCreateCommandQueue(clContext, device_id, 0, &errcode);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in creating command queue");
|
||||
|
||||
// allocate opencl buffers
|
||||
clInBuff = clCreateBuffer(clContext, CL_MEM_READ_WRITE, ELEMENTS * sizeof(uint32_t), NULL, &errcode);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in creating buffer");
|
||||
clOutBuff = clCreateBuffer(clContext, CL_MEM_READ_WRITE, ELEMENTS * sizeof(uint32_t), NULL, &errcode);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in creating buffer");
|
||||
|
||||
// init buffers
|
||||
errcode = clEnqueueWriteBuffer(clCommandQue, clInBuff, CL_TRUE, 0, ELEMENTS * sizeof(uint32_t), m_in, 0, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in writing to buffer");
|
||||
errcode = clEnqueueWriteBuffer(clCommandQue, clOutBuff, CL_TRUE, 0, ELEMENTS * sizeof(uint32_t), (uint32_t*)m_out, 0, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in writing to buffer");
|
||||
|
||||
// create a program from the kernel source
|
||||
const size_t kernel_size = test_Gen9core_gen_len;
|
||||
const unsigned char* kernel_bin = test_Gen9core_gen;
|
||||
clProgram = clCreateProgramWithBinary(clContext, 1, &device_id, &kernel_size, &kernel_bin, NULL, &errcode);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in loading binary");
|
||||
|
||||
// build the program
|
||||
errcode = clBuildProgram(clProgram, 1, &device_id, NULL, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in building program");
|
||||
|
||||
// create the OpenCL kernel
|
||||
clKernel = clCreateKernel(clProgram, "clmain", &errcode);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in creating kernel");
|
||||
|
||||
// set kernel args
|
||||
errcode = clSetKernelArg(clKernel, 0, sizeof(cl_mem), (void *)&clInBuff);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in setting kernel arg");
|
||||
errcode = clSetKernelArg(clKernel, 1, sizeof(cl_mem), (void *)&clOutBuff);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in setting kernel arg");
|
||||
|
||||
// launch the kernel
|
||||
size_t globalWorkSize = ELEMENTS;
|
||||
errcode = clEnqueueNDRangeKernel(clCommandQue, clKernel, 1, NULL, &globalWorkSize, NULL, 0, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in launching kernel");
|
||||
|
||||
// wait for finish
|
||||
clFinish(clCommandQue);
|
||||
|
||||
// read result back
|
||||
errcode = clEnqueueReadBuffer(clCommandQue, clOutBuff, CL_TRUE, 0, ELEMENTS * sizeof(uint32_t), (void*)m_out, 0, NULL, NULL);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in reading GPU mem");
|
||||
|
||||
uint32_t errors = 0;
|
||||
for(int i = 0; i < ELEMENTS; i++)
|
||||
{
|
||||
if(m_out[i] != num)
|
||||
{
|
||||
//LOG_INFO("Error at Item " << i << " val: " << m_out[i]);
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
Genode::log("Task has finished with ", errors, " errors!");
|
||||
|
||||
// free stuff
|
||||
errcode = clReleaseKernel(clKernel);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in releasing kernel");
|
||||
errcode = clReleaseMemObject(clInBuff);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in releasing mem obj");
|
||||
errcode = clReleaseMemObject(clOutBuff);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in releasing mem obj");
|
||||
errcode = clReleaseCommandQueue(clCommandQue);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in releasing command queue");
|
||||
errcode = clReleaseContext(clContext);
|
||||
if(errcode != CL_SUCCESS) Genode::log("Error in releasing context");
|
||||
|
||||
// free buffers
|
||||
alloc.free(m_in);
|
||||
alloc.free((void*)m_out);
|
||||
}
|
||||
12
repos/hello_gpgpu/src/hello_gpgpu/test.h
Normal file
12
repos/hello_gpgpu/src/hello_gpgpu/test.h
Normal file
@@ -0,0 +1,12 @@
|
||||
#ifndef TEST_H
|
||||
#define TEST_H
|
||||
|
||||
#include <base/allocator_avl.h>
|
||||
|
||||
/**
|
||||
* @brief run a test kernel
|
||||
*
|
||||
*/
|
||||
void run_gpgpu_test(Genode::Allocator_avl& alloc);
|
||||
|
||||
#endif // TEST_H
|
||||
44
repos/mml/include/kiihdytin/gpgpu/client.h
Normal file
44
repos/mml/include/kiihdytin/gpgpu/client.h
Normal file
@@ -0,0 +1,44 @@
|
||||
/*
|
||||
* \brief CLient-side interface to a GPGPU session
|
||||
* \author Michael Müller
|
||||
* \date 2022-07-17
|
||||
*/
|
||||
|
||||
/*
|
||||
* Copyright (C) 2022 Michael Müller
|
||||
*
|
||||
* This file is distributed under the terms of the
|
||||
* GNU Affero General Public License version 3.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <base/rpc_client.h>
|
||||
#include <base/log.h>
|
||||
#include "session.h"
|
||||
#include "kernel.h"
|
||||
|
||||
namespace Kiihdytin::GPGPU {
|
||||
struct Session_client;
|
||||
}
|
||||
|
||||
struct Kiihdytin::GPGPU::Session_client : Genode::Rpc_client<Kiihdytin::GPGPU::Session>
|
||||
{
|
||||
Session_client(Genode::Capability<Session> cap)
|
||||
: Genode::Rpc_client<Session>(cap) { }
|
||||
|
||||
void enqueue_kernel(Kernel &kernel) override {
|
||||
call<Rpc_enqueue_kernel>(kernel);
|
||||
}
|
||||
|
||||
void wait_for_kernel(Kernel &kernel) override {
|
||||
call<Rpc_wait_for_kernel>(kernel);
|
||||
}
|
||||
|
||||
void abort_kernel(Kernel &kernel) override {
|
||||
call<Rpc_abort_kernel>(kernel);
|
||||
}
|
||||
|
||||
void remove_kernel(Kernel &kernel) override {
|
||||
call<Rpc_remove_kernel>(kernel);
|
||||
}
|
||||
};
|
||||
30
repos/mml/include/kiihdytin/gpgpu/connection.h
Normal file
30
repos/mml/include/kiihdytin/gpgpu/connection.h
Normal file
@@ -0,0 +1,30 @@
|
||||
/*
|
||||
* \brief Connection to GPGPU session
|
||||
* \author Michael Müller
|
||||
* \date 2022-07-17
|
||||
*/
|
||||
|
||||
/*
|
||||
* Copyright (C) 2022 Michael Müller
|
||||
*
|
||||
* This file is distributed under the terms of the
|
||||
* GNU Affero General Public License version 3.
|
||||
*/
|
||||
|
||||
#include "client.h"
|
||||
#include <base/connection.h>
|
||||
|
||||
namespace Kiihdytin::GPGPU { struct Connection; }
|
||||
|
||||
|
||||
struct Kiihdytin::GPGPU::Connection : Genode::Connection<Session>, Session_client
|
||||
{
|
||||
Connection(Genode::Env &env)
|
||||
:
|
||||
/* create session */
|
||||
Genode::Connection<Kiihdytin::GPGPU::Session>(env, session(env.parent(),
|
||||
"ram_quota=6K, cap_quota=4")), // TODO: determine correct ram and cap quota
|
||||
|
||||
/* initialize RPC interface */
|
||||
Session_client(cap()) { }
|
||||
};
|
||||
46
repos/mml/include/kiihdytin/gpgpu/kernel.h
Normal file
46
repos/mml/include/kiihdytin/gpgpu/kernel.h
Normal file
@@ -0,0 +1,46 @@
|
||||
/*
|
||||
* \brief Definition of a GPGPU kernel, i.e. OpenCL-slang for an executable unit of code for an OpenCL device
|
||||
* \author Michael Müller
|
||||
* \date 2022-07-15
|
||||
*/
|
||||
|
||||
/*
|
||||
* Copyright (C) 2022 Michael Müller
|
||||
*
|
||||
* This file is distributed under the terms of the
|
||||
* GNU Affero General Public License version 3.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <gpgpu_driver.h>
|
||||
#include <cstdint>
|
||||
#include <util/list.h>
|
||||
#include <driver/lib/chain.h>
|
||||
|
||||
namespace Kiihdytin::GPGPU {
|
||||
|
||||
typedef uint8_t Kernel_image;
|
||||
/**
|
||||
* @class This class represents an OpenCL kernel
|
||||
*
|
||||
*/
|
||||
class Kernel : public Chain
|
||||
{
|
||||
private:
|
||||
struct kernel_config _configuration;
|
||||
Kernel_image *_image;
|
||||
|
||||
public:
|
||||
/**
|
||||
* @brief get configuration for this kernel
|
||||
* @return reference to kernel configuration
|
||||
*/
|
||||
inline struct kernel_config &get_config() { return _configuration; }
|
||||
|
||||
/**
|
||||
* @brief get pointer to kernel image
|
||||
* @return pointer to kernel's binary image
|
||||
*/
|
||||
inline Kernel_image *get_image() { return _image; }
|
||||
};
|
||||
}
|
||||
92
repos/mml/include/kiihdytin/gpgpu/scheduler.h
Normal file
92
repos/mml/include/kiihdytin/gpgpu/scheduler.h
Normal file
@@ -0,0 +1,92 @@
|
||||
/*
|
||||
* \brief Scheduler interface for the GPGPU, select which vGPU to choose next.
|
||||
* \author Michael Müller
|
||||
* \date 2022-07-15
|
||||
*/
|
||||
|
||||
/*
|
||||
* Copyright (C) 2022 Michael Müller
|
||||
*
|
||||
* This file is distributed under the terms of the
|
||||
* GNU Affero General Public License version 3.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <driver/lib/wf_queue.h>
|
||||
#include <driver/gpgpu_driver.h>
|
||||
#include <driver/ppgtt32.h>
|
||||
#include "vgpu.h"
|
||||
#include "kernel.h"
|
||||
|
||||
namespace Kiihdytin::GPGPU {
|
||||
|
||||
class Scheduler
|
||||
{
|
||||
private:
|
||||
VGpu *_curr_vgpu;
|
||||
// GPGPU_Driver _driver; /* TODO: Use driver session */
|
||||
WFQueue _run_list;
|
||||
/* TODO: Define handler object for GPGPU driver session to receive interrupts. */
|
||||
|
||||
public:
|
||||
|
||||
Scheduler()
|
||||
{
|
||||
// TODO: Initialize GPU driver
|
||||
|
||||
// TODO: Register interrupt/event handler for the GPGPU driver session.
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Select next vGPU from run list
|
||||
* @details At the moment, round-robin is the only implemented strategy.
|
||||
* TODO: Implement interface for strategies and strategies *
|
||||
*/
|
||||
void schedule_next() {
|
||||
VGpu *next;
|
||||
|
||||
if ((next = static_cast<VGpu*>(_run_list.dequeue()))) {
|
||||
this->dispatch(*next);
|
||||
_curr_vgpu = next;
|
||||
_run_list.enqueue(next);
|
||||
} else
|
||||
_curr_vgpu = nullptr;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Switch to new vGPU's context
|
||||
*
|
||||
* @param vgpu - vGPU to switch to
|
||||
*/
|
||||
void dispatch(VGpu &vgpu) {
|
||||
// TODO: Implement context switch using GPGPU driver
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Implmentation for the handling of events from the GPU
|
||||
* @details The handler is especially important for scheduling the next vGPU and for
|
||||
* executing kernels. It is the target for interrupts coming from the GPGPU driver, e.g. when
|
||||
* a kernel has finished its execution.
|
||||
*/
|
||||
void handle_gpu_event() {
|
||||
// TODO: Check for error conditions
|
||||
|
||||
// TODO: Handle finish of kernel
|
||||
|
||||
/* Switch to next vGPU in the run list */
|
||||
schedule_next();
|
||||
|
||||
/* If no vGPU to schedule, this means that we don't have any clients anymore.
|
||||
* Thus, there are also no kernels anymore to run. */
|
||||
if (_curr_vgpu == nullptr)
|
||||
return;
|
||||
|
||||
Kernel *next = _curr_vgpu->take_kernel();
|
||||
|
||||
if (!next) /* If there is no kernel for the vGPU left */
|
||||
schedule_next(); /* pick the next vGPU, maybe it has got some kernels for us. */
|
||||
|
||||
// TODO: execute kernel using GPGPU driver
|
||||
}
|
||||
};
|
||||
}
|
||||
44
repos/mml/include/kiihdytin/gpgpu/service.h
Normal file
44
repos/mml/include/kiihdytin/gpgpu/service.h
Normal file
@@ -0,0 +1,44 @@
|
||||
/*
|
||||
* \brief Definition of the GPGPU service's root component
|
||||
* \author Michael Müller
|
||||
* \date 2022-07-17
|
||||
*/
|
||||
|
||||
/*
|
||||
* Copyright (C) 2022 Michael Müller
|
||||
*
|
||||
* This file is distributed under the terms of the
|
||||
* GNU Affero General Public License version 3.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <base/component.h>
|
||||
#include <root/component.h>
|
||||
#include <base/heap.h>
|
||||
#include <util/list.h>
|
||||
#include "session_component.h"
|
||||
|
||||
|
||||
namespace Kiihdytin::GPGPU {
|
||||
class Service;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief The GPGPU service provides multiplexed accesses to GPGPU functionality to its clients.
|
||||
*
|
||||
*/
|
||||
|
||||
class Kiihdytin::GPGPU::Session : public Genode::Root_component<Kiihdytin::GPGPU::Session_component>
|
||||
{
|
||||
private:
|
||||
Genode::List<Kiihdytin::GPGPU::Session_component> sessions;
|
||||
|
||||
protected:
|
||||
|
||||
Session_component *_create_session(const char*) override {
|
||||
return new (md_alloc()) Session_component();
|
||||
}
|
||||
|
||||
public:
|
||||
Session(Genode::Entrypoint &ep, Genode::Allocator &alloc) : Genode::Root_component<Session_component>(ep, alloc) {}
|
||||
};
|
||||
54
repos/mml/include/kiihdytin/gpgpu/session.h
Normal file
54
repos/mml/include/kiihdytin/gpgpu/session.h
Normal file
@@ -0,0 +1,54 @@
|
||||
/*
|
||||
* \brief Interface definition of a GPU session
|
||||
* \author Michael Müller
|
||||
* \date 2022-07-17
|
||||
*/
|
||||
|
||||
/*
|
||||
* Copyright (C) 2022 Michael Müller
|
||||
*
|
||||
* This file is distributed under the terms of the
|
||||
* GNU Affero General Public License version 3.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <session/session.h>
|
||||
#include <base/rpc.h>
|
||||
#include "vgpu.h"
|
||||
|
||||
namespace Kiihdytin::GPGPU {
|
||||
class Session;
|
||||
}
|
||||
|
||||
class Kiihdytin::GPGPU::Session : Genode::Session
|
||||
{
|
||||
|
||||
private:
|
||||
VGpu &vgpu;
|
||||
|
||||
VGpu& create_vgpu();
|
||||
PPGTT32& create_ppgtt();
|
||||
|
||||
|
||||
public:
|
||||
static const char *service_name() { return "Kiihdytin::GPGPU"; }
|
||||
|
||||
enum { CAP_QUOTA = 2 }; // TODO: determine actual cap quota
|
||||
|
||||
Session() : vgpu(create_vgpu()) {}
|
||||
|
||||
/* Backend methods */
|
||||
virtual void enqueue_kernel(Kernel &kernel) = 0;
|
||||
virtual void wait_for_kernel(Kernel &kernel) = 0;
|
||||
virtual void abort_kernel(Kernel &kernel) = 0;
|
||||
virtual void remove_kernel(Kernel &kernel) = 0;
|
||||
|
||||
/* RPC interface */
|
||||
|
||||
GENODE_RPC(Rpc_enqueue_kernel, void, enqueue_kernel, Kernel&);
|
||||
GENODE_RPC(Rpc_wait_for_kernel, void, wait_for_kernel, Kernel &);
|
||||
GENODE_RPC(Rpc_abort_kernel, void, abort_kernel, Kernel &);
|
||||
GENODE_RPC(Rpc_remove_kernel, void, remove_kernel, Kernel &);
|
||||
|
||||
GENODE_RPC_INTERFACE(Rpc_enqueue_kernel, Rpc_remove_kernel, Rpc_wait_for_kernel, Rpc_abort_kernel);
|
||||
};
|
||||
31
repos/mml/include/kiihdytin/gpgpu/session_component.h
Normal file
31
repos/mml/include/kiihdytin/gpgpu/session_component.h
Normal file
@@ -0,0 +1,31 @@
|
||||
/*
|
||||
* \brief RPC object for a GPGPU session
|
||||
* \author Michael Müller
|
||||
* \date 2022-07-17
|
||||
*/
|
||||
|
||||
/*
|
||||
* Copyright (C) 2022 Michael Müller
|
||||
*
|
||||
* This file is distributed under the terms of the
|
||||
* GNU Affero General Public License version 3.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <base/component.h>
|
||||
#include <base/rpc_server.h>
|
||||
#include "session.h"
|
||||
|
||||
namespace Kiihdytin::GPGPU {
|
||||
class Session_component;
|
||||
}
|
||||
|
||||
class Kiihdytin::GPGPU::Session_component : public Genode::Rpc_object<Kiihdytin::GPGPU::Session>
|
||||
{
|
||||
public:
|
||||
void enqueue_kernel(Kernel &kernel) override;
|
||||
void wait_for_kernel(Kernel &kernel) override;
|
||||
void abort_kernel(Kernel &kernel) override;
|
||||
void remove_kernel(Kernel &kernel) override;
|
||||
|
||||
};
|
||||
73
repos/mml/include/kiihdytin/gpgpu/vgpu.h
Normal file
73
repos/mml/include/kiihdytin/gpgpu/vgpu.h
Normal file
@@ -0,0 +1,73 @@
|
||||
/*
|
||||
* \brief Representation of a "virtual" GPU, used as abstraction for the real thing.
|
||||
* \author Michael Müller
|
||||
* \date 2022-07-15
|
||||
*/
|
||||
|
||||
/*
|
||||
* Copyright (C) 2022 Michael Müller
|
||||
*
|
||||
* This file is distributed under the terms of the
|
||||
* GNU Affero General Public License version 3.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <driver/gpgpu_driver.h>
|
||||
#include "kernel.h"
|
||||
#include <driver/ppgtt32.h>
|
||||
#include <driver/lib/chain.h>
|
||||
#include <driver/lib/wf_queue.h>
|
||||
|
||||
namespace Kiihdytin::GPGPU {
|
||||
|
||||
class Context;
|
||||
|
||||
class VGpu : public Chain
|
||||
{
|
||||
private:
|
||||
// Context _context; TODO: implement context images
|
||||
PPGTT32 &_ppgtt;
|
||||
WFQueue _ready_list;
|
||||
|
||||
public:
|
||||
/**
|
||||
* @brief Construct a new VGpu object
|
||||
*
|
||||
* @param ppgtt - PPGTT mapping phyisical addresses from the client's rm space to gpu addresses
|
||||
*/
|
||||
VGpu(PPGTT32 &ppgtt) : _ppgtt(ppgtt) {}
|
||||
|
||||
/**
|
||||
* @brief Add a kernel to the vGPU's ready list
|
||||
*
|
||||
* @param kernel - the kernel object to enqueue
|
||||
*/
|
||||
void add_kernel(Kernel &kernel) {
|
||||
_ready_list.enqueue(&kernel);
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Get saved GPU context for this VGPU
|
||||
*
|
||||
* @return GPU context image for this VGPU
|
||||
* TODO: implement saving the context of the GPU using the GPGPU driver
|
||||
*/
|
||||
Context get_context();
|
||||
|
||||
/**
|
||||
* @brief Dequeue a kernel from the ready list
|
||||
*
|
||||
* @return First kernel image in ready list
|
||||
*/
|
||||
Kernel *take_kernel() { return static_cast<Kernel*>(_ready_list.dequeue()); }
|
||||
|
||||
/**
|
||||
* @brief Get the ppgtt object
|
||||
*
|
||||
* @return PPGTT
|
||||
*/
|
||||
PPGTT32 &get_ppgtt() { return _ppgtt; }
|
||||
|
||||
|
||||
};
|
||||
}
|
||||
@@ -1,87 +1,28 @@
|
||||
#include <base/component.h>
|
||||
#include <timer_session/connection.h>
|
||||
#include <base/heap.h>
|
||||
#include <cstdint>
|
||||
#include <memory>
|
||||
#include <thread>
|
||||
#include <iostream>
|
||||
#include <chrono>
|
||||
|
||||
namespace Thread_test {
|
||||
class Tester;
|
||||
class Test_thread;
|
||||
namespace ThreadTest {
|
||||
struct Main;
|
||||
}
|
||||
|
||||
using namespace Genode;
|
||||
|
||||
class Thread_test::Test_thread : public Thread
|
||||
struct ThreadTest::Main
|
||||
{
|
||||
private:
|
||||
Env &_env;
|
||||
uint16_t _id;
|
||||
Timer::Connection _timer{_env};
|
||||
Genode::Env &_env;
|
||||
|
||||
public:
|
||||
List_element<Test_thread> _list_element{this};
|
||||
|
||||
Test_thread(Env &env, uint16_t id, Location const &location)
|
||||
: Thread(env, Name("test_", location.xpos(), "x", location.ypos()), 4 * 4096, location, Weight(), env.cpu()),
|
||||
_env(env),
|
||||
_id(id)
|
||||
{ }
|
||||
|
||||
void entry() override
|
||||
{
|
||||
while(true) {
|
||||
Genode::log("Pong from thread ", _id);
|
||||
auto start = _timer.elapsed_ms();
|
||||
// auto start = std::chrono::steady_clock::now ();
|
||||
_timer.msleep(_id * 1000);
|
||||
auto end = _timer.elapsed_ms();
|
||||
// auto end = std::chrono::steady_clock::now();
|
||||
Genode::log("Thread ", _id, " woke up afer", (end-start), " ms.");
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
class Thread_test::Tester
|
||||
{
|
||||
typedef List<List_element<Thread_test::Test_thread>> Thread_list;
|
||||
|
||||
private:
|
||||
Env &_env;
|
||||
Heap _heap {_env.ram(), _env.rm()};
|
||||
Thread_list _threads{};
|
||||
|
||||
public:
|
||||
Tester(Env &env) : _env(env)
|
||||
void execute()
|
||||
{
|
||||
Affinity::Space space = env.cpu().affinity_space();
|
||||
|
||||
Genode::log("Size of Affinity space is ", space.total());
|
||||
Genode::log("-----------------------------");
|
||||
for (unsigned i = 1; i < space.total(); i++)
|
||||
{
|
||||
Affinity::Location loc = space.location_of_index(i);
|
||||
Genode::log("1: x = ", loc.xpos(), " y = ", loc.ypos());
|
||||
while(true) {
|
||||
std::cout << "Hello world" << std::endl;
|
||||
std::this_thread::sleep_for(std::chrone::seconds(1));
|
||||
}
|
||||
Genode::log("-----------------------------");
|
||||
|
||||
for (unsigned i = 1; i < space.total(); i++)
|
||||
{
|
||||
Affinity::Location location = env.cpu().affinity_space().location_of_index(i);
|
||||
Test_thread *thread = new (_heap) Test_thread(env, (uint16_t)i, location);
|
||||
thread->start();
|
||||
|
||||
_threads.insert(&thread->_list_element);
|
||||
}
|
||||
/* Test, whether unique_ptrs work */
|
||||
//auto unique_thread = std::unique_ptr<Test_thread>(new (_heap) Test_thread(env, 255, env.cpu().affinity_space().location_of_index(0)));
|
||||
//unique_thread->start();
|
||||
}
|
||||
};
|
||||
|
||||
void Component::construct(Genode::Env &env)
|
||||
{
|
||||
env.exec_static_constructors();
|
||||
static Thread_test::Tester tester(env);
|
||||
Genode::log("Thread tester constructed.");
|
||||
}
|
||||
static ThreadTest::Main main(env);
|
||||
std::thread([main]
|
||||
{ main->execute(); });
|
||||
}
|
||||
Reference in New Issue
Block a user