diff options
Diffstat (limited to 'BS/baselines')
| -rw-r--r-- | BS/baselines/cpu/Makefile | 4 | ||||
| -rw-r--r-- | BS/baselines/cpu/README | 9 | ||||
| -rw-r--r-- | BS/baselines/cpu/bs_omp.c | 110 | ||||
| -rwxr-xr-x | BS/baselines/cpu/timer.h | 59 | ||||
| -rw-r--r-- | BS/baselines/gpu/Makefile | 2 | ||||
| -rw-r--r-- | BS/baselines/gpu/README | 9 | ||||
| -rw-r--r-- | BS/baselines/gpu/binary_search.cu | 125 | ||||
| -rw-r--r-- | BS/baselines/gpu/binary_search.h | 18 | ||||
| -rw-r--r-- | BS/baselines/gpu/cpu_lib.py | 21 | ||||
| -rw-r--r-- | BS/baselines/gpu/cu_lib_import.py | 39 | ||||
| -rw-r--r-- | BS/baselines/gpu/run.py | 22 |
11 files changed, 418 insertions, 0 deletions
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)) |