PrIM -- first commit

7 files changed, 236 insertions, 0 deletions

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))