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/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 +++++++ 11 files changed, 418 insertions(+) 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 (limited to 'BS/baselines') 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)) -- cgit v1.2.3