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/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 +++++++ 7 files changed, 236 insertions(+) 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/gpu') 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