From 3de4b495fb176eba9a0eb517a4ce05903cb67acb Mon Sep 17 00:00:00 2001 From: Juan Gomez Luna Date: Wed, 16 Jun 2021 19:46:05 +0200 Subject: PrIM -- first commit --- BS/Makefile | 44 ++++++++ BS/baselines/cpu/Makefile | 4 + BS/baselines/cpu/README | 9 ++ BS/baselines/cpu/bs_omp.c | 110 ++++++++++++++++++ BS/baselines/cpu/timer.h | 59 ++++++++++ BS/baselines/gpu/Makefile | 2 + BS/baselines/gpu/README | 9 ++ BS/baselines/gpu/binary_search.cu | 125 ++++++++++++++++++++ BS/baselines/gpu/binary_search.h | 18 +++ BS/baselines/gpu/cpu_lib.py | 21 ++++ BS/baselines/gpu/cu_lib_import.py | 39 +++++++ BS/baselines/gpu/run.py | 22 ++++ BS/dpu/task.c | 153 +++++++++++++++++++++++++ BS/host/app.c | 232 ++++++++++++++++++++++++++++++++++++++ BS/support/common.h | 49 ++++++++ BS/support/params.h | 52 +++++++++ BS/support/timer.h | 59 ++++++++++ 17 files changed, 1007 insertions(+) create mode 100644 BS/Makefile create mode 100644 BS/baselines/cpu/Makefile create mode 100644 BS/baselines/cpu/README create mode 100644 BS/baselines/cpu/bs_omp.c create mode 100755 BS/baselines/cpu/timer.h create mode 100644 BS/baselines/gpu/Makefile create mode 100644 BS/baselines/gpu/README create mode 100644 BS/baselines/gpu/binary_search.cu create mode 100644 BS/baselines/gpu/binary_search.h create mode 100644 BS/baselines/gpu/cpu_lib.py create mode 100644 BS/baselines/gpu/cu_lib_import.py create mode 100644 BS/baselines/gpu/run.py create mode 100644 BS/dpu/task.c create mode 100644 BS/host/app.c create mode 100755 BS/support/common.h create mode 100644 BS/support/params.h create mode 100755 BS/support/timer.h (limited to 'BS') 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 +#include +#include +#include +#include +#include +#include +#include +#include +#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 + +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 +#include +#include "binary_search.h" + +#include +#include + +#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<<>>(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(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 + #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 +#include +#include +#include +#include +#include +#include +#include +#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 +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#if ENERGY +#include +#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 # of untimed warmup iterations (default=1)" + "\n -e # of timed repetition iterations (default=3)" + "\n" + "\nBenchmark-specific options:" + "\n -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 + +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)); } -- cgit v1.2.3