PrIM -- first commit

17 files changed, 1007 insertions, 0 deletions

diff --git a/BS/Makefile b/BS/Makefile
new file mode 100644
index 0000000..129ade1
--- /dev/null
+++ b/BS/Makefile
@@ -0,0 +1,44 @@
+DPU_DIR := dpu
+HOST_DIR := host
+BUILDDIR ?= bin
+NR_TASKLETS ?= 16
+NR_DPUS ?= 1
+PROBLEM_SIZE ?= 2
+
+define conf_filename
+	${BUILDDIR}/.NR_DPUS_$(1)_NR_TASKLETS_$(2).conf
+endef
+CONF := $(call conf_filename,${NR_DPUS},${NR_TASKLETS})
+
+COMMON_INCLUDES := support
+HOST_TARGET := ${BUILDDIR}/bs_host
+DPU_TARGET := ${BUILDDIR}/bs_dpu
+
+HOST_SOURCES := $(wildcard ${HOST_DIR}/*.c)
+DPU_SOURCES := $(wildcard ${DPU_DIR}/*.c)
+
+.PHONY: all clean test
+
+__dirs := $(shell mkdir -p ${BUILDDIR})
+
+COMMON_FLAGS := -Wall -Wextra  -g -I${COMMON_INCLUDES}
+HOST_FLAGS := ${COMMON_FLAGS} -std=c11 -O3 `dpu-pkg-config --cflags --libs dpu` -DNR_TASKLETS=${NR_TASKLETS} -DNR_DPUS=${NR_DPUS} -DPROBLEM_SIZE=${PROBLEM_SIZE}
+DPU_FLAGS := ${COMMON_FLAGS} -O2 -DNR_TASKLETS=${NR_TASKLETS}
+
+all: ${HOST_TARGET} ${DPU_TARGET}
+
+${CONF}:
+	$(RM) $(call conf_filename,*,*)
+	touch ${CONF}
+
+${HOST_TARGET}: ${HOST_SOURCES} ${COMMON_INCLUDES} ${CONF}
+	$(CC) -o $@ ${HOST_SOURCES} ${HOST_FLAGS}
+
+${DPU_TARGET}: ${DPU_SOURCES} ${COMMON_INCLUDES} ${CONF}
+	dpu-upmem-dpurte-clang ${DPU_FLAGS} -o $@ ${DPU_SOURCES}
+
+clean:
+	$(RM) -r $(BUILDDIR)
+
+test: all
+	./${HOST_TARGET} -i 262144
diff --git a/BS/baselines/cpu/Makefile b/BS/baselines/cpu/Makefile
new file mode 100644
index 0000000..e907fc8
--- /dev/null
+++ b/BS/baselines/cpu/Makefile
@@ -0,0 +1,4 @@
+all:
+	gcc bs_omp.c -o bs_omp -fopenmp
+run:
+	./bs_omp 262144 16777216
diff --git a/BS/baselines/cpu/README b/BS/baselines/cpu/README
new file mode 100644
index 0000000..ce82830
--- /dev/null
+++ b/BS/baselines/cpu/README
@@ -0,0 +1,9 @@
+Binary Search (BS)
+
+Compilation instructions:
+
+    make
+
+Execution instructions
+
+    ./bs_omp 2048576 16777216
diff --git a/BS/baselines/cpu/bs_omp.c b/BS/baselines/cpu/bs_omp.c
new file mode 100644
index 0000000..3ad83ed
--- /dev/null
+++ b/BS/baselines/cpu/bs_omp.c
@@ -0,0 +1,110 @@
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdbool.h>
+#include <string.h>
+#include <unistd.h>
+#include <getopt.h>
+#include <assert.h>
+#include <time.h>
+#include <stdint.h>
+#include "timer.h"
+
+#define DTYPE uint64_t
+/*
+* @brief creates a "test file" by filling a bufferwith values
+*/
+void create_test_file(DTYPE * input, uint64_t  nr_elements, DTYPE * querys, uint64_t n_querys) {
+
+  uint64_t max = UINT64_MAX;
+  uint64_t min = 0;
+
+  srand(time(NULL));
+
+  input[0] = 1;
+  for (uint64_t i = 1; i < nr_elements; i++) {
+        input[i] = input[i - 1] + (rand() % 10) + 1;
+  }
+
+  for(uint64_t i = 0; i < n_querys; i++)
+  {
+	querys[i] = input[rand() % (nr_elements - 2)];
+  }
+}
+
+/**
+* @brief compute output in the host
+*/
+uint64_t binarySearch(DTYPE * input, uint64_t input_size, DTYPE* querys, unsigned n_querys)
+{
+
+	uint64_t found = -1;
+	uint64_t q, r, l, m;
+	
+       #pragma omp parallel for private(q,r,l,m)
+     	for(q = 0; q < n_querys; q++)
+      	{
+		l = 0;
+		r = input_size;
+		while (l <= r) 
+		{
+	    		m = l + (r - l) / 2;
+
+	    		// Check if x is present at mid
+	     		if (input[m] == querys[q])
+			{	
+		    		found += m;
+				break;
+			}
+	    		// If x greater, ignore left half
+	    		if (input[m] < querys[q])
+			    	l = m + 1;
+
+	    		// If x is smaller, ignore right half
+			else
+		    		r = m - 1;
+		
+		}
+       	}
+
+      	return found;
+}
+
+  /**
+  * @brief Main of the Host Application.
+  */
+  int main(int argc, char **argv) {
+
+    Timer timer;
+    uint64_t input_size = atol(argv[1]);
+    uint64_t n_querys = atol(argv[2]);
+
+    printf("Vector size: %lu, num searches: %lu\n", input_size, n_querys);
+	
+    DTYPE * input = malloc((input_size) * sizeof(DTYPE));
+    DTYPE * querys = malloc((n_querys) * sizeof(DTYPE));
+
+    DTYPE result_host = -1;
+
+    // Create an input file with arbitrary data.
+    create_test_file(input, input_size, querys, n_querys);
+	
+    start(&timer, 0, 0);
+    result_host = binarySearch(input, input_size - 1, querys, n_querys);   
+    stop(&timer, 0);
+
+
+    int status = (result_host);
+    if (status) {
+        printf("[OK] Execution time: ");
+	print(&timer, 0, 1);
+	printf("ms.\n");
+    } else {
+        printf("[ERROR]\n");
+    }
+    free(input);
+
+
+    return status ? 0 : 1;
+}
+
diff --git a/BS/baselines/cpu/timer.h b/BS/baselines/cpu/timer.h
new file mode 100755
index 0000000..969ef97
--- /dev/null
+++ b/BS/baselines/cpu/timer.h
@@ -0,0 +1,59 @@
+/*

+ * Copyright (c) 2016 University of Cordoba and University of Illinois

+ * All rights reserved.

+ *

+ * Developed by:    IMPACT Research Group

+ *                  University of Cordoba and University of Illinois

+ *                  http://impact.crhc.illinois.edu/

+ *

+ * Permission is hereby granted, free of charge, to any person obtaining a copy

+ * of this software and associated documentation files (the "Software"), to deal

+ * with the Software without restriction, including without limitation the 

+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or

+ * sell copies of the Software, and to permit persons to whom the Software is

+ * furnished to do so, subject to the following conditions:

+ *

+ *      > Redistributions of source code must retain the above copyright notice,

+ *        this list of conditions and the following disclaimers.

+ *      > Redistributions in binary form must reproduce the above copyright

+ *        notice, this list of conditions and the following disclaimers in the

+ *        documentation and/or other materials provided with the distribution.

+ *      > Neither the names of IMPACT Research Group, University of Cordoba, 

+ *        University of Illinois nor the names of its contributors may be used 

+ *        to endorse or promote products derived from this Software without 

+ *        specific prior written permission.

+ *

+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR

+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,

+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 

+ * CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER

+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,

+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH

+ * THE SOFTWARE.

+ *

+ */

+

+#include <sys/time.h>

+

+typedef struct Timer{

+

+    struct timeval startTime[4];

+    struct timeval stopTime[4];

+    double         time[4];

+

+}Timer;

+

+void start(Timer *timer, int i, int rep) {

+    if(rep == 0) {

+        timer->time[i] = 0.0;

+    }

+    gettimeofday(&timer->startTime[i], NULL);

+}

+

+void stop(Timer *timer, int i) {

+    gettimeofday(&timer->stopTime[i], NULL);

+    timer->time[i] += (timer->stopTime[i].tv_sec - timer->startTime[i].tv_sec) * 1000000.0 +

+                      (timer->stopTime[i].tv_usec - timer->startTime[i].tv_usec);

+}

+

+void print(Timer *timer, int i, int REP) { printf("%f\t", timer->time[i] / (1000 * REP)); }

diff --git a/BS/baselines/gpu/Makefile b/BS/baselines/gpu/Makefile
new file mode 100644
index 0000000..c5bafab
--- /dev/null
+++ b/BS/baselines/gpu/Makefile
@@ -0,0 +1,2 @@
+all:
+	nvcc -arch=sm_30 -m64 -Xcompiler -fPIC -shared -o cu_binary_search.so binary_search.cu -std=c++11
diff --git a/BS/baselines/gpu/README b/BS/baselines/gpu/README
new file mode 100644
index 0000000..a3d8a0a
--- /dev/null
+++ b/BS/baselines/gpu/README
@@ -0,0 +1,9 @@
+Binary Search (BS)
+
+Compilation instructions:
+
+    make
+
+Execution instructions
+
+    python3 run.py
diff --git a/BS/baselines/gpu/binary_search.cu b/BS/baselines/gpu/binary_search.cu
new file mode 100644
index 0000000..2cb3cb7
--- /dev/null
+++ b/BS/baselines/gpu/binary_search.cu
@@ -0,0 +1,125 @@
+#include <cuda.h>

+#include <limits.h>

+#include "binary_search.h"

+

+#include <chrono>

+#include <iostream>

+

+#define BLOCKDIM 512

+#define SEARCH_CHUNK 16

+#define BLOCK_CHUNK (BLOCKDIM*SEARCH_CHUNK)

+

+

+__global__ void search_kernel(const long int *arr,

+    const long int len, const long int *querys, const long int num_querys, long int *res, bool *flag)

+{

+    int search;

+    if(*flag == false) {

+        int tid = threadIdx.x;

+        __shared__ int s_arr[BLOCK_CHUNK];

+

+        /* Since each value is being copied to shared memory, the rest of the

+        following uncommented code is unncessary, since a direct comparison

+        can be done at the time of copy below. */

+        // for(int i = 0; i < BLOCKDIM; ++i) {

+        //     int shared_loc = i*SEARCH_CHUNK + tid;

+        //     int global_loc = shared_loc + BLOCK_CHUNK * blockIdx.x;

+        //     if(arr[global_loc] == search) {

+        //         *flag = true;

+        //         *res = global_loc;

+        //     }

+        //     __syncthreads();

+        // }

+

+        /* Copy chunk of array that this entire block of threads will read

+        from the slower global memory to the faster shared memory. */

+        for(long int i = 0; i < SEARCH_CHUNK; ++i) {

+            int shared_loc = tid*SEARCH_CHUNK + i;

+            int global_loc = shared_loc + BLOCK_CHUNK * blockIdx.x;

+

+            /* Make sure to stay within the bounds of the global array,

+            else assign a dummy value. */

+            if(global_loc < len) {

+              s_arr[shared_loc] = arr[global_loc];

+            }

+            else {

+              s_arr[shared_loc] = INT_MAX;

+            }

+        }

+        __syncthreads();

+

+        for(long int i = 0; i < num_querys; i++)

+        {

+            search = querys[i];

+            /* For each thread, set the initial search range. */

+            int L = 0;

+            int R = SEARCH_CHUNK - 1;

+            int m = (L + R) / 2;

+

+            /* Pointer to the part of the shared array for this thread. */

+            int *s_ptr = &s_arr[tid*SEARCH_CHUNK];

+

+            /* Each thread will search a chunk of the block array.

+            Many blocks will not find a solution so the search must

+            be allowed to fail on a per block basis. The loop will

+            break (fail) when L >= R. */

+            while(L <= R && *flag == false)

+            {

+                if(s_ptr[m] < search) {

+                    L = m + 1;

+                }

+                else if(s_ptr[m] > search) {

+                    R = m - 1;

+                }

+                else {

+                    *flag = true;

+                    *res = m += tid*SEARCH_CHUNK + BLOCK_CHUNK * blockIdx.x;

+                }

+

+                m = (L + R) / 2;

+            }

+        }

+    }

+}

+

+

+

+int binary_search(const long int *arr, const long int len, const long int *querys, const long int num_querys)

+{

+    long int *d_arr, *d_querys, *d_res;

+    bool *d_flag;

+

+    size_t arr_size = len * sizeof(long int);

+    size_t querys_size = num_querys * sizeof(long int);

+    size_t res_size = sizeof(long int);

+    size_t flag_size = sizeof(bool);

+

+    cudaMalloc(&d_arr, arr_size);

+    cudaMalloc(&d_querys, querys_size);

+    cudaMalloc(&d_res, res_size);

+    cudaMalloc(&d_flag, flag_size);

+

+    cudaMemcpy(d_arr, arr, arr_size, cudaMemcpyHostToDevice);

+    cudaMemcpy(d_querys, querys, querys_size, cudaMemcpyHostToDevice);

+    cudaMemset(d_flag, 0, flag_size);

+

+    /* Set res value to -1, so that if the function returns -1, that

+    indicates an algorithm failure. */

+    cudaMemset(d_res, -0x1, res_size);

+

+    int blockSize = BLOCKDIM;

+    int gridSize = (len-1)/BLOCK_CHUNK + 1;

+

+    auto start = std::chrono::high_resolution_clock::now();

+    search_kernel<<<gridSize,blockSize>>>(d_arr, len, d_querys, num_querys ,d_res, d_flag);

+    cudaDeviceSynchronize();

+    auto end = std::chrono::high_resolution_clock::now();

+    std::cout << "Kernel Time: " <<

+        std::chrono::duration_cast<std::chrono::milliseconds>(end-start).count() <<

+        " ms" << std::endl;

+

+    long int res;

+    cudaMemcpy(&res, d_res, res_size, cudaMemcpyDeviceToHost);

+

+    return res;

+}

diff --git a/BS/baselines/gpu/binary_search.h b/BS/baselines/gpu/binary_search.h
new file mode 100644
index 0000000..5849506
--- /dev/null
+++ b/BS/baselines/gpu/binary_search.h
@@ -0,0 +1,18 @@
+#ifndef BINARY_SEARCH_H

+#define BINARY_SEARCH_H

+

+#ifdef _WIN32

+  #include <windows.h>

+  #define DLL_EXPORT __declspec(dllexport)

+#else

+  #define DLL_EXPORT

+#endif

+

+

+extern "C" {

+

+    int DLL_EXPORT binary_search(const long int *arr, const long int len, const long int *querys, const long int num_querys);

+

+}

+

+#endif /* BINARY_SEARCH_H */

diff --git a/BS/baselines/gpu/cpu_lib.py b/BS/baselines/gpu/cpu_lib.py
new file mode 100644
index 0000000..9a45f94
--- /dev/null
+++ b/BS/baselines/gpu/cpu_lib.py
@@ -0,0 +1,21 @@
+# -*- coding: utf-8 -*-
+
+def binary_search(arr, search):
+    
+    L = 0
+    R = len(arr)
+    
+    while(L<=R):
+        
+        if L>R:
+            return -1 #Error code 1
+        
+        m = (L+R)/2
+        if(arr[m] < search):
+            L = m+1
+        elif(arr[m] > search):
+            R = m-1
+        else:
+            return m
+    
+    return -2 #Error code 2
\ No newline at end of file
diff --git a/BS/baselines/gpu/cu_lib_import.py b/BS/baselines/gpu/cu_lib_import.py
new file mode 100644
index 0000000..aafbbce
--- /dev/null
+++ b/BS/baselines/gpu/cu_lib_import.py
@@ -0,0 +1,39 @@
+# -*- coding: utf-8 -*-

+

+__all__ = [

+    "binary_search",

+]

+

+

+from ctypes import *

+import os.path as path

+from numpy.ctypeslib import load_library, ndpointer

+import platform

+

+

+## Load the DLL

+if platform.system() == 'Linux':

+    cuda_lib = load_library("cu_binary_search.so", path.dirname(path.realpath(__file__)))

+elif platform.system() == 'Windows':

+    cuda_lib = load_library("cu_binary_search.dll", path.dirname(path.realpath(__file__)))

+

+

+

+

+## Define argtypes for all functions to import

+argtype_defs = {

+

+    "binary_search" : [ndpointer("i8"),

+                       c_int,

+                       ndpointer("i8"),

+                       c_int],

+

+}

+

+

+

+

+## Import functions from DLL

+for func, argtypes in argtype_defs.items():

+    locals().update({func: cuda_lib[func]})

+    locals()[func].argtypes = argtypes

diff --git a/BS/baselines/gpu/run.py b/BS/baselines/gpu/run.py
new file mode 100644
index 0000000..58963b9
--- /dev/null
+++ b/BS/baselines/gpu/run.py
@@ -0,0 +1,22 @@
+# -*- coding: utf-8 -*-
+
+import numpy as np
+import time
+
+#Local Imports
+from cu_lib_import import binary_search as gpu_search
+
+# Set an array size to create
+arr_len = 2048576
+num_querys = 16777216
+
+# Dummy array created
+arr = np.arange(0, arr_len, 1).astype("i8")
+
+# Random search querys created
+querys = np.random.randint(1, arr_len, num_querys)
+
+# GPU search function call
+t0 = time.time()
+res_gpu = gpu_search(arr, len(arr), querys, len(querys))
+print("Total GPU Time: %i ms" % ((time.time() - t0)*1e003))
diff --git a/BS/dpu/task.c b/BS/dpu/task.c
new file mode 100644
index 0000000..39a340d
--- /dev/null
+++ b/BS/dpu/task.c
@@ -0,0 +1,153 @@
+/*
+* Binary Search with multiple tasklets
+*
+*/
+#include <stdint.h>
+#include <stdio.h>
+#include <defs.h>
+#include <mram.h>
+#include <alloc.h>
+#include <mram.h>
+#include <barrier.h>
+#include <perfcounter.h>
+#include "common.h"
+
+__host dpu_arguments_t DPU_INPUT_ARGUMENTS;
+__host dpu_results_t DPU_RESULTS[NR_TASKLETS];
+
+// Search
+static DTYPE search(DTYPE *bufferA, DTYPE searching_for) {
+  DTYPE found = -2;
+  if(bufferA[0] <= searching_for)
+  {
+    found = -1;
+    for (uint32_t i = 0; i < BLOCK_SIZE / sizeof(DTYPE); i++){
+      if(bufferA[i] == searching_for)
+      {
+        found = i;
+        break;
+      }
+    }
+  }
+  return found;
+}
+
+BARRIER_INIT(my_barrier, NR_TASKLETS);
+
+extern int main_kernel1(void);
+
+int(*kernels[nr_kernels])(void) = {main_kernel1};
+
+int main(void){
+  // Kernel
+  return kernels[DPU_INPUT_ARGUMENTS.kernel]();
+}
+
+// main_kernel1
+int main_kernel1() {
+  unsigned int tasklet_id = me();
+  #if PRINT
+  printf("tasklet_id = %u\n", tasklet_id);
+  #endif
+  if(tasklet_id == 0){
+    mem_reset(); // Reset the heap
+  }
+  // Barrier
+  barrier_wait(&my_barrier);
+
+  DTYPE searching_for, found;
+  uint64_t input_size = DPU_INPUT_ARGUMENTS.input_size;
+
+  // Address of the current processing block in MRAM
+  uint32_t start_mram_block_addr_A       = (uint32_t) DPU_MRAM_HEAP_POINTER;
+  uint32_t start_mram_block_addr_aux     = start_mram_block_addr_A;
+  uint32_t end_mram_block_addr_A         = start_mram_block_addr_A + sizeof(DTYPE) * input_size;
+  uint32_t current_mram_block_addr_query = end_mram_block_addr_A + tasklet_id * (DPU_INPUT_ARGUMENTS.slice_per_dpu / NR_TASKLETS) * sizeof(DTYPE);
+
+  // Initialize a local cache to store the MRAM block
+  DTYPE *cache_A     = (DTYPE *) mem_alloc(BLOCK_SIZE);
+  DTYPE *cache_aux_A = (DTYPE *) mem_alloc(BLOCK_SIZE);
+  DTYPE *cache_aux_B = (DTYPE *) mem_alloc(BLOCK_SIZE);
+
+  dpu_results_t *result = &DPU_RESULTS[tasklet_id];
+
+  for(uint64_t targets = 0; targets < (DPU_INPUT_ARGUMENTS.slice_per_dpu / NR_TASKLETS); targets++)
+  {
+    found = -1;
+
+    mram_read((__mram_ptr void const *) current_mram_block_addr_query, &searching_for, 8);
+    current_mram_block_addr_query += 8;
+
+    bool end = false;
+
+    // Initialize input vector boundaries
+    start_mram_block_addr_A    = (uint32_t) DPU_MRAM_HEAP_POINTER;
+    start_mram_block_addr_aux  = start_mram_block_addr_A;
+    end_mram_block_addr_A      = start_mram_block_addr_A + sizeof(DTYPE) * input_size;
+
+    uint32_t current_mram_block_addr_A = start_mram_block_addr_A;
+
+    // Bring first and last values to WRAM
+    mram_read((__mram_ptr void const *) current_mram_block_addr_A, cache_aux_A, BLOCK_SIZE);
+    mram_read((__mram_ptr void const *) (end_mram_block_addr_A - BLOCK_SIZE * sizeof(DTYPE)),   cache_aux_B, BLOCK_SIZE);
+
+    current_mram_block_addr_A = (start_mram_block_addr_A + end_mram_block_addr_A) / 2;
+    while(!end)
+    {
+      // Load cache with current MRAM block
+      mram_read((__mram_ptr void const *) current_mram_block_addr_A, cache_A, BLOCK_SIZE);
+
+      // Search inside block
+      found = search(cache_A, searching_for);
+
+      // If found > -1, we found the searching_for query
+      if(found > -1)
+      {
+        result->found = found + (current_mram_block_addr_A - start_mram_block_addr_aux) / sizeof(DTYPE);
+        break;
+      }
+
+      // If found == -2, we need to discard right part of the input vector
+      if(found == -2)
+      {
+        end_mram_block_addr_A     = current_mram_block_addr_A;
+        current_mram_block_addr_A = (current_mram_block_addr_A + start_mram_block_addr_A) / 2;
+      }
+
+      // If found == -1, we need to discard left part of the input vector
+      else if (found == -1)
+      {
+        start_mram_block_addr_A   = current_mram_block_addr_A;
+        current_mram_block_addr_A = (current_mram_block_addr_A + end_mram_block_addr_A) / 2;
+      }
+
+      // Start boundary check
+      if(current_mram_block_addr_A < (start_mram_block_addr_aux + BLOCK_SIZE))
+      {
+        end = true;
+        mram_read((__mram_ptr void const *) current_mram_block_addr_A, cache_A, BLOCK_SIZE);
+        found = search(cache_A, searching_for);
+
+        if(found > -1)
+        {
+          end = true;
+          result->found = found + (current_mram_block_addr_A - start_mram_block_addr_aux) / sizeof(DTYPE);
+        }
+      }
+
+      // End boundary check
+      if(current_mram_block_addr_A > (end_mram_block_addr_A - BLOCK_SIZE))
+      {
+        end = true;
+        mram_read((__mram_ptr void const *) end_mram_block_addr_A - BLOCK_SIZE, cache_A, BLOCK_SIZE);
+        found = search(cache_A, searching_for);
+
+        if(found > -1)
+        {
+          result->found = found + (current_mram_block_addr_A - start_mram_block_addr_aux) / sizeof(DTYPE);
+        }
+      }
+    }
+  }
+  return 0;
+}
diff --git a/BS/host/app.c b/BS/host/app.c
new file mode 100644
index 0000000..3929877
--- /dev/null
+++ b/BS/host/app.c
@@ -0,0 +1,232 @@
+/**
+* app.c
+* BS Host Application Source File
+*
+*/
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdbool.h>
+#include <string.h>
+#include <dpu.h>
+#include <dpu_log.h>
+#include <unistd.h>
+#include <getopt.h>
+#include <assert.h>
+#include <time.h>
+
+#if ENERGY
+#include <dpu_probe.h>
+#endif
+
+#include "params.h"
+#include "timer.h"
+
+// Define the DPU Binary path as DPU_BINARY here
+#define DPU_BINARY "./bin/bs_dpu"
+
+// Create input arrays
+void create_test_file(DTYPE * input, DTYPE * querys, uint64_t  nr_elements, uint64_t nr_querys) {
+
+	input[0] = 1;
+	for (uint64_t i = 1; i < nr_elements; i++) {
+		input[i] = input[i - 1] + 1;
+	}
+	for (uint64_t i = 0; i < nr_querys; i++) {
+		querys[i] = i;
+	}
+}
+
+// Compute output in the host
+int64_t binarySearch(DTYPE * input, DTYPE * querys, DTYPE input_size, uint64_t num_querys)
+{
+	uint64_t result = -1;
+	DTYPE r;
+	for(uint64_t q = 0; q < num_querys; q++)
+	{
+		DTYPE l = 0;
+		r = input_size;
+		while (l <= r) {
+			DTYPE m = l + (r - l) / 2;
+
+			// Check if x is present at mid
+			if (input[m] == querys[q])
+			result = m;
+
+			// If x greater, ignore left half
+			if (input[m] < querys[q])
+			l = m + 1;
+
+			// If x is smaller, ignore right half
+			else
+			r = m - 1;
+		}
+	}
+	return result;
+}
+
+
+// Main of the Host Application
+int main(int argc, char **argv) {
+
+	struct Params p = input_params(argc, argv);
+	struct dpu_set_t dpu_set, dpu;
+	uint32_t nr_of_dpus;
+	uint64_t input_size = INPUT_SIZE;
+	uint64_t num_querys = p.num_querys;
+	DTYPE result_host = -1;
+	DTYPE result_dpu  = -1;
+
+	// Create the timer
+	Timer timer;
+
+	// Allocate DPUs and load binary
+	DPU_ASSERT(dpu_alloc(NR_DPUS, NULL, &dpu_set));
+	DPU_ASSERT(dpu_load(dpu_set, DPU_BINARY, NULL));
+	DPU_ASSERT(dpu_get_nr_dpus(dpu_set, &nr_of_dpus));
+
+	#if ENERGY
+	struct dpu_probe_t probe;
+	DPU_ASSERT(dpu_probe_init("energy_probe", &probe));
+	#endif
+
+	// Query number adjustement for proper partitioning
+	if(num_querys % (nr_of_dpus * NR_TASKLETS))
+	num_querys = num_querys + (nr_of_dpus * NR_TASKLETS - num_querys % (nr_of_dpus * NR_TASKLETS));
+
+	assert(num_querys % (nr_of_dpus * NR_TASKLETS) == 0 && "Input dimension");    // Allocate input and querys vectors
+
+	DTYPE * input  = malloc((input_size) * sizeof(DTYPE));
+	DTYPE * querys = malloc((num_querys) * sizeof(DTYPE));
+
+	// Create an input file with arbitrary data
+	create_test_file(input, querys, input_size, num_querys);
+
+	// Compute host solution
+	start(&timer, 0, 0);
+	result_host = binarySearch(input, querys, input_size - 1, num_querys);
+	stop(&timer, 0);
+
+	// Create kernel arguments
+	uint64_t slice_per_dpu          = num_querys / nr_of_dpus;
+	dpu_arguments_t input_arguments = {input_size, slice_per_dpu, 0};
+
+	for (unsigned int rep = 0; rep < p.n_warmup + p.n_reps; rep++) {
+		// Perform input transfers
+		uint64_t i = 0;
+
+		if (rep >= p.n_warmup)
+		start(&timer, 1, rep - p.n_warmup);
+
+		DPU_FOREACH(dpu_set, dpu, i)
+		{
+			DPU_ASSERT(dpu_prepare_xfer(dpu, &input_arguments));
+		}
+
+		DPU_ASSERT(dpu_push_xfer(dpu_set, DPU_XFER_TO_DPU, "DPU_INPUT_ARGUMENTS", 0, sizeof(input_arguments), DPU_XFER_DEFAULT));
+
+		i = 0;
+
+		DPU_FOREACH(dpu_set, dpu, i)
+		{
+			DPU_ASSERT(dpu_prepare_xfer(dpu, input));
+		}
+
+		DPU_ASSERT(dpu_push_xfer(dpu_set, DPU_XFER_TO_DPU, DPU_MRAM_HEAP_POINTER_NAME, 0, input_size * sizeof(DTYPE), DPU_XFER_DEFAULT));
+
+		i = 0;
+
+		DPU_FOREACH(dpu_set, dpu, i)
+		{
+			DPU_ASSERT(dpu_prepare_xfer(dpu, querys + slice_per_dpu * i));
+		}
+
+		DPU_ASSERT(dpu_push_xfer(dpu_set, DPU_XFER_TO_DPU, DPU_MRAM_HEAP_POINTER_NAME, input_size * sizeof(DTYPE), slice_per_dpu * sizeof(DTYPE), DPU_XFER_DEFAULT));
+
+		if (rep >= p.n_warmup)
+		stop(&timer, 1);
+
+		// Run kernel on DPUs
+		if (rep >= p.n_warmup)
+		{
+			start(&timer, 2, rep - p.n_warmup);
+			#if ENERGY
+			DPU_ASSERT(dpu_probe_start(&probe));
+			#endif
+		}
+
+		DPU_ASSERT(dpu_launch(dpu_set, DPU_SYNCHRONOUS));
+
+		if (rep >= p.n_warmup)
+		{
+			stop(&timer, 2);
+			#if ENERGY
+			DPU_ASSERT(dpu_probe_stop(&probe));
+			#endif
+		}
+		// Print logs if required
+		#if PRINT
+		unsigned int each_dpu = 0;
+		printf("Display DPU Logs\n");
+		DPU_FOREACH(dpu_set, dpu)
+		{
+			printf("DPU#%d:\n", each_dpu);
+			DPU_ASSERT(dpulog_read_for_dpu(dpu.dpu, stdout));
+			each_dpu++;
+		}
+		#endif
+
+		// Retrieve results
+		if (rep >= p.n_warmup)
+		start(&timer, 3, rep - p.n_warmup);
+		dpu_results_t* results_retrieve[nr_of_dpus];
+		i = 0;
+		DPU_FOREACH(dpu_set, dpu, i)
+		{
+			results_retrieve[i] = (dpu_results_t*)malloc(NR_TASKLETS * sizeof(dpu_results_t));
+			DPU_ASSERT(dpu_prepare_xfer(dpu, results_retrieve[i]));
+		}
+
+		DPU_ASSERT(dpu_push_xfer(dpu_set, DPU_XFER_FROM_DPU, "DPU_RESULTS", 0, NR_TASKLETS * sizeof(dpu_results_t), DPU_XFER_DEFAULT));
+
+		DPU_FOREACH(dpu_set, dpu, i)
+		{
+			for(unsigned int each_tasklet = 0; each_tasklet < NR_TASKLETS; each_tasklet++)
+			{
+				if(results_retrieve[i][each_tasklet].found > result_dpu)
+				{
+					result_dpu = results_retrieve[i][each_tasklet].found;
+				}
+			}
+			free(results_retrieve[i]);
+		}
+		if(rep >= p.n_warmup)
+		stop(&timer, 3);
+	}
+	// Print timing results
+	printf("CPU Version Time (ms): ");
+	print(&timer, 0, p.n_reps);
+	printf("CPU-DPU Time (ms): ");
+	print(&timer, 1, p.n_reps);
+	printf("DPU Kernel Time (ms): ");
+	print(&timer, 2, p.n_reps);
+	printf("DPU-CPU Time (ms): ");
+	print(&timer, 3, p.n_reps);
+
+	#if ENERGY
+	double energy;
+	DPU_ASSERT(dpu_probe_get(&probe, DPU_ENERGY, DPU_AVERAGE, &energy));
+	printf("DPU Energy (J): %f\t", energy * num_iterations);
+	#endif
+
+	int status = (result_dpu == result_host);
+	if (status) {
+		printf("[" ANSI_COLOR_GREEN "OK" ANSI_COLOR_RESET "] results are equal\n");
+	} else {
+		printf("[" ANSI_COLOR_RED "ERROR" ANSI_COLOR_RESET "] results differ!\n");
+	}
+
+	free(input);
+	DPU_ASSERT(dpu_free(dpu_set));
+
+	return status ? 0 : 1;
+}
diff --git a/BS/support/common.h b/BS/support/common.h
new file mode 100755
index 0000000..413a6f8
--- /dev/null
+++ b/BS/support/common.h
@@ -0,0 +1,49 @@
+#ifndef _COMMON_H_
+#define _COMMON_H_
+
+#ifdef TL
+#define TASKLETS_INITIALIZER TASKLETS(TL, main, 2048, 2)
+#define NB_OF_TASKLETS_PER_DPU TL
+#else
+#define TASKLETS_INITIALIZER TASKLETS(16, main, 2048, 2)
+#define NB_OF_TASKLETS_PER_DPU 16
+#endif
+
+// Transfer size between MRAM and WRAM
+#ifdef BL
+#define BLOCK_SIZE_LOG2 BL
+#define BLOCK_SIZE (1 << BLOCK_SIZE_LOG2)
+#else
+#define BLOCK_SIZE_LOG2 8
+#define BLOCK_SIZE (1 << BLOCK_SIZE_LOG2)
+#endif
+
+// Data type
+#define DTYPE int64_t
+
+// Vector size
+#define INPUT_SIZE 2048576
+
+typedef struct {
+	uint64_t input_size;
+	uint64_t slice_per_dpu;
+	enum kernels {
+		kernel1 = 0,
+		nr_kernels = 1,
+	} kernel;
+} dpu_arguments_t;
+
+// Structures used by both the host and the dpu to communicate information
+typedef struct {
+    DTYPE found;
+} dpu_results_t;
+
+#ifndef ENERGY
+#define ENERGY 0
+#endif
+#define PRINT 0
+
+#define ANSI_COLOR_RED     "\x1b[31m"
+#define ANSI_COLOR_GREEN   "\x1b[32m"
+#define ANSI_COLOR_RESET   "\x1b[0m"
+#endif
diff --git a/BS/support/params.h b/BS/support/params.h
new file mode 100644
index 0000000..02bd750
--- /dev/null
+++ b/BS/support/params.h
@@ -0,0 +1,52 @@
+#ifndef _PARAMS_H_
+#define _PARAMS_H_
+
+#include "common.h"
+
+typedef struct Params {
+  long  num_querys;
+  unsigned   n_warmup;
+  unsigned   n_reps;
+}Params;
+
+void usage() {
+  fprintf(stderr,
+    "\nUsage:  ./program [options]"
+    "\n"
+    "\nGeneral options:"
+    "\n    -h        help"
+    "\n    -w <W>    # of untimed warmup iterations (default=1)"
+    "\n    -e <E>    # of timed repetition iterations (default=3)"
+    "\n"
+    "\nBenchmark-specific options:"
+    "\n    -i <I>    problem size (default=2 queries)"
+    "\n");
+  }
+
+  struct Params input_params(int argc, char **argv) {
+    struct Params p;
+    p.num_querys    = PROBLEM_SIZE;
+    p.n_warmup      = 1;
+    p.n_reps        = 3;
+
+    int opt;
+    while((opt = getopt(argc, argv, "h:i:w:e:")) >= 0) {
+      switch(opt) {
+        case 'h':
+        usage();
+        exit(0);
+        break;
+        case 'i': p.num_querys    = atol(optarg); break;
+        case 'w': p.n_warmup      = atoi(optarg); break;
+        case 'e': p.n_reps        = atoi(optarg); break; 
+	default:
+        	fprintf(stderr, "\nUnrecognized option!\n");
+        	usage();
+        	exit(0);
+      }
+    }
+    assert(NR_DPUS > 0 && "Invalid # of dpus!");
+
+    return p;
+  }
+  #endif
diff --git a/BS/support/timer.h b/BS/support/timer.h
new file mode 100755
index 0000000..969ef97
--- /dev/null
+++ b/BS/support/timer.h
@@ -0,0 +1,59 @@
+/*

+ * Copyright (c) 2016 University of Cordoba and University of Illinois

+ * All rights reserved.

+ *

+ * Developed by:    IMPACT Research Group

+ *                  University of Cordoba and University of Illinois

+ *                  http://impact.crhc.illinois.edu/

+ *

+ * Permission is hereby granted, free of charge, to any person obtaining a copy

+ * of this software and associated documentation files (the "Software"), to deal

+ * with the Software without restriction, including without limitation the 

+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or

+ * sell copies of the Software, and to permit persons to whom the Software is

+ * furnished to do so, subject to the following conditions:

+ *

+ *      > Redistributions of source code must retain the above copyright notice,

+ *        this list of conditions and the following disclaimers.

+ *      > Redistributions in binary form must reproduce the above copyright

+ *        notice, this list of conditions and the following disclaimers in the

+ *        documentation and/or other materials provided with the distribution.

+ *      > Neither the names of IMPACT Research Group, University of Cordoba, 

+ *        University of Illinois nor the names of its contributors may be used 

+ *        to endorse or promote products derived from this Software without 

+ *        specific prior written permission.

+ *

+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR

+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,

+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 

+ * CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER

+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,

+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH

+ * THE SOFTWARE.

+ *

+ */

+

+#include <sys/time.h>

+

+typedef struct Timer{

+

+    struct timeval startTime[4];

+    struct timeval stopTime[4];

+    double         time[4];

+

+}Timer;

+

+void start(Timer *timer, int i, int rep) {

+    if(rep == 0) {

+        timer->time[i] = 0.0;

+    }

+    gettimeofday(&timer->startTime[i], NULL);

+}

+

+void stop(Timer *timer, int i) {

+    gettimeofday(&timer->stopTime[i], NULL);

+    timer->time[i] += (timer->stopTime[i].tv_sec - timer->startTime[i].tv_sec) * 1000000.0 +

+                      (timer->stopTime[i].tv_usec - timer->startTime[i].tv_usec);

+}

+

+void print(Timer *timer, int i, int REP) { printf("%f\t", timer->time[i] / (1000 * REP)); }