From 14e6ffcc19a47e79325fcf18778c443473ed775b Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Mon, 19 Jun 2023 02:05:56 +0800 Subject: feat(exp): Add a argument in test Signed-off-by: KunoiSayami --- CMakeLists.txt | 1 - analysis0617.py | 20 +- expt_0525.cu | 19 +- expt_0528.cu | 66 ++--- expt_0618.cu | 788 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++ expt_0618_3.cu | 10 +- 6 files changed, 836 insertions(+), 68 deletions(-) create mode 100644 expt_0618.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index bef2d9d..f2e65e0 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -176,7 +176,6 @@ set_cuda_target_base(expt_0529) add_executable(expt_0604 expt_0604.cpp) set_target_properties(expt_0604 PROPERTIES LINKER_LANGUAGE CXX) - add_executable(expt_0618 expt_0618.cu) set_cuda_target_base(expt_0618) diff --git a/analysis0617.py b/analysis0617.py index 46a36b9..2e1fe7a 100755 --- a/analysis0617.py +++ b/analysis0617.py @@ -8,21 +8,22 @@ import random async def grab_output(p: asyncio.subprocess.Process) -> float | None: out = (await p.communicate())[0].decode() if p.returncode: - print("Error:", out) + print("Error:", p.returncode, out.strip()) return None print(out.strip()) return float(out.rsplit(':')[-1]) -async def run_custom_exec(p_s: str, sample: int, search: int, insertion: int) -> float | None: - p = (await asyncio.create_subprocess_exec(p_s, str(sample), str(search), str(insertion), - stdout=asyncio.subprocess.PIPE)) +async def run_custom_exec(p_s: str, sample: int, search: int, insertion: int, total_row: str) -> float | None: + p = (await asyncio.create_subprocess_exec(p_s, str(sample), str(search), str(insertion), total_row, + stdout=asyncio.subprocess.PIPE, stderr=asyncio.subprocess.STDOUT)) await p.wait() return await grab_output(p) -async def run_exec(p_s: str, search: int, insertion: int) -> float | None: - p = await asyncio.create_subprocess_exec(p_s, str(search), str(insertion), stdout=asyncio.subprocess.PIPE) +async def run_exec(p_s: str, search: int, insertion: int, total_row: str) -> float | None: + p = await asyncio.create_subprocess_exec(p_s, str(search), str(insertion), total_row, + stdout=asyncio.subprocess.PIPE, stderr=asyncio.subprocess.STDOUT) await p.wait() return await grab_output(p) @@ -78,7 +79,7 @@ async def main(matches: argparse.Namespace) -> None: for _ in range(3): await asyncio.sleep(1) retries = 3 - while (normal := await run_exec(matches.exec2, search, insertion)) is None: + while (normal := await run_exec(matches.exec2, search, insertion, matches.total_row)) is None: retries -= 1 fault_run += 1 if not retries: @@ -86,7 +87,9 @@ async def main(matches: argparse.Namespace) -> None: total_run += 4 - retries await asyncio.sleep(1) retries = 3 - while (custom := await run_custom_exec(matches.exec1, sample, search, insertion)) is None: + while ( + custom := await run_custom_exec(matches.exec1, sample, search, insertion, + matches.total_row)) is None: retries -= 1 fault_run += 1 if not retries: @@ -108,6 +111,7 @@ if __name__ == '__main__': run_.add_argument("exec2") run_.add_argument("sample_limit") run_.add_argument("limit") + run_.add_argument("--total-row", default=0) run_.add_argument("--test-times", default=3) parser_ = arg_.parse_args() if parser_.sub == 'run': diff --git a/expt_0525.cu b/expt_0525.cu index 8e7cd10..6450282 100644 --- a/expt_0525.cu +++ b/expt_0525.cu @@ -596,14 +596,19 @@ inline auto calcBlocks(size_t input) { } int main(int argc, char **argv) { - if (argc != 4) { - printf("Usage %s [sample] [search] [insertion]\n", argv[0]); + if (argc < 4) { + printf("Usage %s [file length]\n", argv[0]); exit(1); } auto sample_length = strtol(argv[1], nullptr, 10); auto search_length = strtol(argv[2], nullptr, 10); auto insertion_length = strtol(argv[3], nullptr, 10); + auto total_row = 0L; + + if (argc > 4) { + total_row = strtol(argv[4], nullptr, 10); + } if (insertion_length < search_length) { printf("Search should smaller than insertion\n"); @@ -611,10 +616,16 @@ int main(int argc, char **argv) { printf("Sample: %ld, Insertion: %ld, Search: %ld ", sample_length, insertion_length, search_length); - fflush(stdout); + ReadHelper readHelper("normal_distribution.txt", sample_length, - insertion_length, 0); + insertion_length, total_row); + + if (total_row) { + printf("Skip: %ld ", readHelper.random_number); + fflush(stdout); + } + readHelper.readFile(); std::vector _sample, _population; readHelper.split_into(_sample, _population); diff --git a/expt_0528.cu b/expt_0528.cu index 27bc29f..bf60070 100644 --- a/expt_0528.cu +++ b/expt_0528.cu @@ -94,7 +94,7 @@ constexpr size_t MAX_LEVEL = 32; // Number of threads per block // #define NUM_THREADS 512 -constexpr size_t NUM_THREADS = 512; +constexpr size_t NUM_THREADS = 2; // constexpr size_t NUM_ITEMS = BUILD_SIZE; // constexpr size_t KEYS = 1048576; @@ -211,52 +211,6 @@ public: } }; -/*struct MapNode { - LL key; - // Node *point[MAX_LEVEL + 1]; - Node *point; -}; - -class MemMap { -public: - size_t size; - size_t real_size; - MapNode *store; - - MemMap() : size(0), store(nullptr), real_size(0) {} - - __device__ bool insert(MapNode node) { - bool need_extend = this->size + 1 > this->real_size; - if (need_extend) { - bool need_copy = this->real_size == 0; - if (!need_copy) { - this->real_size += 1; - } - this->real_size *= 2; - MapNode *old = this->store; - this->store = new MapNode[this->real_size]; - if (need_copy) { - memcpy(this->store, old, this->real_size * sizeof(MapNode *)); - } - delete[] old; - } - // need sort after insert - this->store[size] = node; - this->size += 1; - } - - __device__ Node *search(LL key) { - for (int offset = 0; offset < this->size; offset++) { - if (this->store[offset].key >= key) { - return this->store[offset].point; - } - } - return nullptr; - } - - __device__ ~MemMap() { delete[] store; } -};*/ - // Definition of lock-free skip list class LockFreeSkipList { @@ -637,22 +591,32 @@ int *generateRandomLevel(size_t size) { } int main(int argc, char **argv) { - if (argc != 3) { - printf("Usage %s [sample] [search] [insertion]\n", argv[0]); + if (argc < 3) { + printf("Usage %s [total row]\n", argv[0]); exit(1); } auto search_length = strtol(argv[1], nullptr, 10); auto insertion_length = strtol(argv[2], nullptr, 10); + auto total_row = 0L; + + if (argc > 3) { + total_row = strtol(argv[3], nullptr, 10); + } if (insertion_length < search_length) { printf("Search should smaller than insertion\n"); } printf("Insertion: %ld, Search: %ld ", insertion_length, search_length); - fflush(stdout); - ReadHelper readHelper("normal_distribution.txt", 0, insertion_length, 0); + + ReadHelper readHelper("normal_distribution.txt", 0, insertion_length, + total_row); + if (total_row) { + printf("Skip: %ld ", readHelper.random_number); + fflush(stdout); + } readHelper.readFile(); std::vector _search(readHelper.population_vector.begin(), readHelper.population_vector.begin() + diff --git a/expt_0618.cu b/expt_0618.cu new file mode 100644 index 0000000..37d2a05 --- /dev/null +++ b/expt_0618.cu @@ -0,0 +1,788 @@ +/* + +Copyright 2012-2013 Indian Institute of Technology Kanpur. All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + +1. Redistributions of source code must retain the above copyright notice, this +list of conditions, and the following disclaimer. +2. Redistributions in binary form must reproduce the above copyright notice, +this list of conditions, and the following disclaimer in the documentation +and/or other materials provided with the distribution. + +THIS SOFTWARE IS PROVIDED BY INDIAN INSTITUTE OF TECHNOLOGY KANPUR ``AS IS'' +AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +DISCLAIMED. IN NO EVENT SHALL INDIAN INSTITUTE OF TECHNOLOGY KANPUR OR +THE CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, +EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT +OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE +POSSIBILITY OF SUCH DAMAGE. + +The views and conclusions contained in the software and documentation are +those of the authors and should not be interpreted as representing official +policies, either expressed or implied, of Indian Institute of Technology Kanpur. + +*/ + +/********************************************************************************** + + Lock-free skip list for CUDA; tested for CUDA 4.2 on 32-bit Ubuntu 10.10 and +64-bit Ubuntu 12.04. Developed at IIT Kanpur. + + Inputs: Percentage of add and delete operations (e.g., 30 50 for 30% add and +50% delete) Output: Prints the total time (in milliseconds) to execute the the +sequence of operations + + Compilation flags: -O3 -arch sm_20 -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc/ +-DNUM_ITEMS=num_ops -DFACTOR=num_ops_per_thread -DKEYS=num_keys + + NUM_ITEMS is the total number of operations (mix of add, delete, search) to +execute. + + FACTOR is the number of operations per thread. + + KEYS is the number of integer keys assumed in the range [10, 9+KEYS]. + The paper cited below states that the key range is [0, KEYS-1]. However, we +have shifted the range by +10 so that the head sentinel key (the minimum key) +can be chosen as zero. Any positive shift other than +10 would also work. + + The include path ~/NVIDIA_GPU_Computing_SDK/C/common/inc/ is needed for +cutil.h. + + Related work: + + Prabhakar Misra and Mainak Chaudhuri. Performance Evaluation of Concurrent +Lock-free Data Structures on GPUs. In Proceedings of the 18th IEEE International +Conference on Parallel and Distributed Systems, December 2012. + +***************************************************************************************/ + +// #include"cutil.h" // Comment this if cutil.h is not available +// #include "cuda_runtime.h" +#include "sortlib.cuh" +#include +#include +#include +#include +#include + +typedef unsigned long long LL; + +// #define MEASURE_TIME +// #define MEASURE_ACCESS + +#include "read_helper.h" + +#if (defined(MEASURE_ACCESS) && defined(MEASURE_TIME)) +#error "Shouldn't define MEASURE_TIME and MEASURE_ACCESS at the same time" +#endif + +#ifdef MEASURE_TIME +#undef BUILD_SIZE +#define BUILD_SIZE 1024 +#endif + +// Maximum level of a node in the skip list +// #define MAX_LEVEL 32 +constexpr size_t MAX_LEVEL = 25; + +// Number of threads per block +// #define NUM_THREADS 512 +constexpr size_t NUM_THREADS = 256; + +// constexpr size_t NUM_ITEMS = BUILD_SIZE; +// constexpr size_t KEYS = 1048576; +constexpr size_t FACTOR = 1; + +// should change this to dynamic next time +// constexpr size_t KEY_INDEX_SIZE = 32; +// constexpr size_t SAMPLE_SIZE = 1023; + +// constexpr int block_size = STEP_SIZE; + +typedef LL key_type; + +#define CUDA_ERROR_CHECK + +#define CudaSafeCall(err) __cudaSafeCall(err, __FILE__, __LINE__) +#define CudaCheckError() __cudaCheckError(__FILE__, __LINE__) + +inline void cudaSafeCall_(cudaError err, const char *file, const int line) { +#ifdef CUDA_ERROR_CHECK + if (cudaSuccess != err) { + fprintf(stderr, "cudaSafeCall() failed at %s:%i : %s\n", file, line, + cudaGetErrorString(err)); + exit(-1); + } +#endif +} + +inline void __cudaCheckError(const char *file, const int line) { +#ifdef CUDA_ERROR_CHECK + cudaError err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "cudaCheckError() failed at %s:%i : %s\n", file, line, + cudaGetErrorString(err)); + exit(-1); + } + + // More careful checking. However, this will affect performance. + // Comment away if needed. + err = cudaDeviceSynchronize(); + if (cudaSuccess != err) { + fprintf(stderr, "cudaCheckError() with sync failed at %s:%i : %s\n", file, + line, cudaGetErrorString(err)); + exit(-1); + } +#endif +} + +// class Node; + +// Definition of generic node class + +class +#ifndef _MSC_VER + __attribute__((aligned(16))) +#else + __declspec(align(16)) +#endif + Node { +public: + int topLevel; // Level of the node + LL key; // Key value + LL next[MAX_LEVEL + 1]{}; // Array of next links + + // Create a next field from a reference and mark bit + static __device__ __host__ LL CreateRef(Node *ref, bool mark) { + LL val = (LL)ref; + val = val | mark; + return val; + } + + __device__ __host__ void SetRef(int index, Node *ref, bool mark) { + next[index] = CreateRef(ref, mark); + } + + // Extract the reference from a next field + __device__ Node *GetReference(int index) { + LL ref = next[index]; + return (Node *)((ref >> 1) << 1); + } + + // Extract the reference and mark bit from a next field + __device__ Node *Get(int index, bool *marked) { + marked[0] = next[index] % 2; + return (Node *)((next[index] >> 1) << 1); + } + + // CompareAndSet wrapper + __device__ bool CompareAndSet(int index, Node *expectedRef, Node *newRef, + bool oldMark, bool newMark) { + LL oldVal = (LL)expectedRef | oldMark; + LL newVal = (LL)newRef | newMark; + LL *ref = &(next[index]); + LL oldValOut = atomicCAS(ref, oldVal, newVal); + if (oldValOut == oldVal) + return true; + return false; + } + + // Constructor for sentinel nodes + explicit Node(LL k) { + key = k; + topLevel = MAX_LEVEL; + int i; + for (i = 0; i < MAX_LEVEL + 1; i++) { + next[i] = CreateRef((Node *)nullptr, false); + } + } +}; + +// Definition of lock-free skip list + +class LockFreeSkipList { + + key_type *sample = nullptr; + size_t sampleLength; + + CustomSort customSort; + + double scaleSize; + +public: + Node *head = nullptr; + Node *tail = nullptr; + int *levels = nullptr; + LockFreeSkipList(key_type *_sample, size_t sample_length, double scale_size) + : sampleLength(sample_length), + customSort(sample_length, sizeof(key_type) * 8), scaleSize(scale_size) { + Node *h = new Node(0); + // size_ = 0; + Node *t = new Node(std::numeric_limits::max() - 1); + cudaMalloc(&head, sizeof(Node)); + + cudaMalloc(&tail, sizeof(Node)); + for (auto i = 0; i < h->topLevel + 1; i++) { + h->SetRef(i, tail, false); + } + cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice); + + cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice); + + cudaMalloc(&this->sample, sizeof(key_type) * sampleLength); + cudaMemcpy(this->sample, _sample, sizeof(key_type) * sampleLength, + cudaMemcpyHostToDevice); + + cudaMalloc(&this->levels, sizeof(int) * 35); + size_t local_level[35]{0}; + cudaMemcpy(this->levels, local_level, sizeof(int) * 8, + cudaMemcpyHostToDevice); + } + __device__ bool find(LL, Node **, Node **); // Helping method + __device__ bool Add(LL); + __device__ bool Delete(LL); + __device__ bool Search(LL); + //~LockFreeSkipList(){cudaFree()} + + __device__ size_t searchIndex(key_type key) { + auto result = customSort.binary_search(this->sample, key); + auto index = customSort.calculate_index(result - this->sample); + return index; + } + + static __device__ unsigned trailing_zeroes(size_t index) { + constexpr auto block_size = 2; + unsigned bits = 0; + LL x = index / block_size; + + if (x) { + while (x % block_size == 0) { + ++bits; + x /= block_size; + } + } + return bits; + } + + __device__ unsigned calcIndex(key_type key) { + auto result = customSort.sample_cdf_custom_version(this->sample, key); + auto cdf_index = result / scaleSize; + auto index = trailing_zeroes((size_t)cdf_index); + return index; + } + +#ifdef MEASURE_ACCESS + unsigned access_times = 0; + + __device__ unsigned getAccessCount() const { return this->access_times; } + __device__ void increaseAccessCount(unsigned count = 1) { + atomicAdd(&this->access_times, count); + } +#else + __device__ void increaseAccessCount(unsigned = 1) {} +#endif + +#ifdef MEASURE_TIME + unsigned round = 0; + __device__ void increaseRoundCount(unsigned count = 1) { + atomicAdd(&this->round, count); + } + int spend_time[NUM_ITEMS]{0}; + unsigned long long total_time = 0; + __device__ unsigned getRoundCount() const { return this->round; } +#endif +}; + +__device__ Node **nodes; // Pool of pre-allocated nodes +__device__ unsigned int pointerIndex = 0; // Index into pool of free nodes +//__device__ LL *randoms; // Array storing the levels of the nodes in the free +// pool +__device__ unsigned int NODE_LIMIT; + +// Function for creating a new node when requested by an add operation + +__device__ Node *GetNewNode(LL key, size_t topLevel) { + LL ind = atomicInc(&pointerIndex, NODE_LIMIT); + Node *n = nodes[ind]; + n->key = key; + // n->topLevel = randoms[ind]; + n->topLevel = (int)topLevel; + int i; + for (i = 0; i < n->topLevel + 1; i++) { + n->SetRef(i, nullptr, false); + } + return n; +} + +__device__ LockFreeSkipList *lockFreeSkipList; // The lock-free skip list + +//__device__ LL KeyIndex[KEY_INDEX_SIZE]; + +//__device__ key_type SampleStorage[SAMPLE_SIZE]; + +// Kernel for initializing device memory + +__global__ void init(LockFreeSkipList *l1, Node **n, + unsigned int insertion_limit) { + // randoms = rands; + nodes = n; + lockFreeSkipList = l1; + NODE_LIMIT = insertion_limit; +} + +// Find the window holding key +// On the way clean up logically deleted nodes (those with set marked bit) + +__device__ bool +LockFreeSkipList::find(LL key, Node **preds, + Node **succs) { // preds and succs are arrays of pointers + int bottomLevel = 0; + bool marked[] = {false}; + bool snip; + Node *pred; + Node *curr; + Node *succ; + bool beenThereDoneThat; + while (true) { + beenThereDoneThat = false; + pred = head; + int level; + for (level = MAX_LEVEL; level >= bottomLevel; level--) { + curr = pred->GetReference(level); + while (true) { + succ = curr->Get(level, marked); + while (marked[0]) { + snip = pred->CompareAndSet(level, curr, succ, false, false); + beenThereDoneThat = true; + if (!snip) + break; + curr = pred->GetReference(level); + succ = curr->Get(level, marked); + beenThereDoneThat = false; + // printf("find key is %d \n",(int)key); + } + if (beenThereDoneThat) + break; + if (curr->key <= key) { + pred = curr; + curr = succ; + } else { + break; + } + } + if (beenThereDoneThat) + break; + preds[level] = pred; + succs[level] = curr; + } + if (beenThereDoneThat) + continue; + return ((curr->key == key)); + } +} + +__device__ bool LockFreeSkipList::Search(LL key) { + int bottomLevel = 0; + bool marked = false; + Node *pred = head; + Node *curr = nullptr; + Node *succ; + int level; + for (level = MAX_LEVEL; level >= bottomLevel; level--) { + curr = pred->GetReference(level); +#ifdef MEASURE_ACCESS + this->increaseAccessCount(); +#endif + while (true) { + succ = curr->Get(level, &marked); +#ifdef MEASURE_ACCESS + this->increaseAccessCount(); +#endif + while (marked) { + curr = curr->GetReference(level); + succ = curr->Get(level, &marked); +#ifdef MEASURE_ACCESS + this->increaseAccessCount(2); +#endif + } + if (curr->key < key) { + pred = curr; + curr = succ; + } else { + break; + } + } + } + return (curr != nullptr && curr->key == key); +} + +__device__ bool LockFreeSkipList::Delete(LL key) { + int bottomLevel = 0; + Node *preds[MAX_LEVEL + 1]; + Node *succs[MAX_LEVEL + 1]; + Node *succ; + bool marked[] = {false}; + while (true) { + bool found = find(key, preds, succs); + if (!found) { + return false; + } else { + Node *nodeToDelete = succs[bottomLevel]; + int level; + for (level = nodeToDelete->topLevel; level >= bottomLevel + 1; level--) { + succ = nodeToDelete->Get(level, marked); + while (!marked[0]) { + nodeToDelete->CompareAndSet(level, succ, succ, false, true); + succ = nodeToDelete->Get(level, marked); + } + } + succ = nodeToDelete->Get(bottomLevel, marked); + while (true) { + bool iMarkedIt = + nodeToDelete->CompareAndSet(bottomLevel, succ, succ, false, true); + succ = succs[bottomLevel]->Get(bottomLevel, marked); + if (iMarkedIt) { + find(key, preds, succs); + // size_ -= 1; + // atomicDec(&size_, 1); + return true; + } else if (marked[0]) { + return false; + } + } + } + } +} + +__device__ bool LockFreeSkipList::Add(LL key) { + auto top_level = calcIndex(key); + atomicAdd(this->levels + top_level, 1); + + Node *newNode = GetNewNode(key, top_level); + int topLevel = newNode->topLevel; + int bottomLevel = 0; + Node *preds[MAX_LEVEL + 1]; + Node *succs[MAX_LEVEL + 1]; + int level; + while (true) { + bool found = find(key, preds, succs); + if (found) { + return false; + } else { + Node *pred; + Node *succ; + for (level = bottomLevel; level <= topLevel; level++) { + succ = succs[level]; + newNode->SetRef(level, succ, false); + } + pred = preds[bottomLevel]; + succ = succs[bottomLevel]; + bool t; + // printf("--- key is %d pred is %d succ is %d level is %d + // \n",(int)key,(int)pred->key,(int)succ->key,0); + t = pred->CompareAndSet(bottomLevel, succ, newNode, false, false); + if (!t) { + continue; + } + for (level = bottomLevel + 1; level <= topLevel; level++) { + + while (true) { + pred = preds[level]; + succ = succs[level]; + newNode->SetRef(level, succ, false); + // printf("-- key is %d pred is %d succ is %d level is %d + // \n",(int)key,(int)pred->key,(int)succ->key,(int)level); + if (pred->CompareAndSet(level, succ, newNode, false, false)) { + break; + } + // printf("key is %d pred is %d succ is %d level is %d + // \n",(int)key,(int)pred->key,(int)succ->key,(int)level); + find(key, preds, succs); + } + } + // size_ += 1; + // this->key_map.insert(MapNode(ll, newNode)); + // atomicAdd(&size_, 1); + return true; + } + } +} + +__global__ void print_level() { + for (int i = 0; i < MAX_LEVEL; i++) { + if (!lockFreeSkipList->levels[i]) { + continue; + } + printf("%d: %d\n", i, lockFreeSkipList->levels[i]); + } +} + +__global__ void print() { + // For debugging + int tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid == 0) { + Node *p = lockFreeSkipList->head; + bool marked = false; + while (p != nullptr) { +#if __WORDSIZE == 64 + printf("%#llx, %u, marked=%u, address is %p : ", p->key, p->topLevel, + marked, p); +#else + printf("%#x, %u, marked=%u, address is %p\n", p->key, p->topLevel, marked, + p); +#endif + for (int i = 0; i < p->topLevel + 1; i++) { + printf(" %d ", (int)(p->GetReference(i)->key)); + } + printf("\n"); + p = p->Get(0, &marked); + } + printf("\n"); + } +} + +// The main kernel + +__global__ void kernel(const LL *items, size_t search_length, LL *result) { + // The array items holds the sequence of keys + // The array op holds the sequence of operations + // The array result, at the end, will hold the outcome of the operations + + for (int i = 0; i < FACTOR; + i++) { // FACTOR is the number of operations per thread + auto tid = + i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x; + if (tid >= search_length) + return; + + // Grab the operation and the associated key and execute + LL item = items[tid]; +#ifdef MEASURE_TIME + unsigned long long start_time = clock64(); +#endif + result[tid] = lockFreeSkipList->Search(item); + assert(result[tid]); +#ifdef MEASURE_TIME + unsigned long long end_time = clock64() - start_time; + if (lockFreeSkipList->spend_time[tid]) { + printf("conflict: %d\n", tid); + } + lockFreeSkipList->spend_time[tid] = (int)end_time; +#endif + } +} + +__global__ void kernelAdd(LL *item, size_t insertion_length) { + + for (int i = 0; i < FACTOR; + i++) { // FACTOR is the number of operations per thread + auto tid = + i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x; + + if (tid >= insertion_length) + return; + lockFreeSkipList->Add(item[tid]); + } +} + +// Generate the level of a newly created node + +__global__ void print_function() { +#ifdef MEASURE_ACCESS + printf("count: %u\n", l->getAccessCount()); +#endif +} + +/*__global__ void copy_function(int *spend_time) { + memcpy(spend_time, lockFreeSkipList->spend_time, sizeof(int) * NUM_ITEMS); +}*/ + +inline double calcSliceSize(size_t insertion_size) { + return 1.0 / (double)insertion_size; +} + +inline auto calcBlocks(size_t input) { + return (input % (NUM_THREADS * FACTOR) == 0) + ? input / (NUM_THREADS * FACTOR) + : (input / (NUM_THREADS * FACTOR)) + 1; +} + +int main(int argc, char **argv) { + if (argc != 4) { + printf("Usage %s [sample] [search] [insertion]\n", argv[0]); + exit(1); + } + + auto sample_length = strtol(argv[1], nullptr, 10); + auto search_length = strtol(argv[2], nullptr, 10); + auto insertion_length = strtol(argv[3], nullptr, 10); + + if (insertion_length < search_length) { + printf("Search should smaller than insertion\n"); + } + + printf("Sample: %ld, Insertion: %ld, Search: %ld ", sample_length, + insertion_length, search_length); + + fflush(stdout); + ReadHelper readHelper("normal_distribution.txt", sample_length, + insertion_length, 0); + readHelper.readFile(); + std::vector _sample, _population; + readHelper.split_into(_sample, _population); + /*printf("%lu, Create search vector: %ld\n", _population.size(), + _population.end() - (_population.begin() + insertion_length));*/ + std::vector _search(_population.begin(), + _population.begin() + search_length); + //_population = _sample; + /*std::vector _tmp_sample = _sample; + _population = _tmp_sample;*/ + rebuildSort(_sample); + //_population.resize(insertion_length); + + // Allocate necessary arrays + // LL *op = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); + // LL *levels = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); + // LL *items = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); + LL *result = new LL[search_length]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); + + // Allocate device memory + + LL *cudaOperatorItems; + // LL *Cop; + LL *cudaResult; + + cudaMalloc(&cudaResult, sizeof(key_type) * search_length); + cudaMalloc(&cudaOperatorItems, sizeof(key_type) * insertion_length); + // cudaMalloc(&Cop, sizeof(LL) * NUM_ITEMS); + // cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS, + // cudaMemcpyHostToDevice); + cudaMemcpy(cudaOperatorItems, _population.data(), + sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice); + // cudaMemcpy(Cop, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice); + Node **pointers = + (Node **)new LL[insertion_length]; // malloc(sizeof(LL) * adds); + Node **Cpointers; + + // Allocate the pool of free nodes + + for (int i = 0; i < insertion_length; i++) { + cudaMalloc(&pointers[i], sizeof(Node)); + } + cudaMalloc(&Cpointers, sizeof(Node *) * insertion_length); + cudaMemcpy(Cpointers, pointers, sizeof(Node *) * insertion_length, + cudaMemcpyHostToDevice); + + // Allocate the skip list + + LockFreeSkipList *Clist; + auto *list = new LockFreeSkipList(_sample.data(), _sample.size(), + calcSliceSize(insertion_length)); + + cudaMalloc(&Clist, sizeof(LockFreeSkipList)); + cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice); + // Calculate the number of thread blocks + // NUM_ITEMS = total number of operations to execute + // NUM_THREADS = number of threads per block + // FACTOR = number of operations per thread + + size_t blocks = calcBlocks(insertion_length); + + // Initialize the device memory + init<<<1, 32>>>(Clist, Cpointers, insertion_length); + cudaDeviceSynchronize(); + + // Insertion to skiplist + kernelAdd<<>>(cudaOperatorItems, insertion_length); + CudaCheckError(); + // cudaDeviceSynchronize(); + + // Re-allocate memory for search + cudaFree(cudaOperatorItems); + cudaMalloc(&cudaOperatorItems, sizeof(key_type) * search_length); + cudaMemcpy(cudaOperatorItems, _search.data(), + sizeof(key_type) * search_length, cudaMemcpyHostToDevice); + + // Launch main kernel + blocks = calcBlocks(search_length); + + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start, nullptr); + + kernel<<>>(cudaOperatorItems, search_length, cudaResult); + CudaCheckError(); + cudaDeviceSynchronize(); + cudaEventRecord(stop, nullptr); + cudaEventSynchronize(stop); + float time; + cudaEventElapsedTime(&time, start, stop); + cudaEventDestroy(start); + cudaEventDestroy(stop); + + // Print kernel execution time in milliseconds + + printf("%lu: %lf\n", search_length, time); + // Check for errors + + // Move results back to host memory + + cudaMemcpy(result, cudaResult, sizeof(LL) * search_length, + cudaMemcpyDeviceToHost); + + // Uncomment the following for debugging + // print<<<1,32>>>(); + cudaDeviceSynchronize(); + + print_level<<<1, 1>>>(); + cudaDeviceSynchronize(); + +#if (defined(MEASURE_TIME) || defined(MEASURE_ACCESS)) + + print_function<<<1, 1>>>(); + + cudaDeviceSynchronize(); +#ifdef MEASURE_TIME + { + int *cuda_tmp = nullptr; + cudaMalloc(&cuda_tmp, sizeof(int) * NUM_ITEMS); + copy_function<<<1, 1>>>(cuda_tmp); + cudaDeviceSynchronize(); + CudaCheckError(); + FILE *file = fopen("spend_time.txt", "w"); + int *tmp = new int[NUM_ITEMS]; + cudaMemcpy(tmp, cuda_tmp, sizeof(int) * NUM_ITEMS, cudaMemcpyDeviceToHost); + // memcpy(tmp, SpendTime, sizeof(int) * NUM_ITEMS); + for (int i = 0; i < NUM_ITEMS; i++) { + if (tmp[i] == 0) + break; + fprintf(file, "%d\n", tmp[i]); + } + // printf("%d\n", i); + delete[] tmp; + fclose(file); + } + // for (auto element : SpendTimeVec) + // printf("%d\n", element); +#endif +#endif + + cudaFree(Clist); + cudaFree(cudaResult); + cudaFree(cudaOperatorItems); + cudaFree(Cpointers); + for (int i = 0; i < insertion_length; i++) { + cudaFree(pointers[i]); + } + cudaFree(pointers); + delete[] result; + delete list; + return 0; +} diff --git a/expt_0618_3.cu b/expt_0618_3.cu index 4558f54..15af8f1 100644 --- a/expt_0618_3.cu +++ b/expt_0618_3.cu @@ -84,7 +84,8 @@ int main(int argc, char const *argv[]) { printf("sample length: %ld, population length: %ld\n", sample_length, population_length); - ReadHelper readHelper("normal_distribution.txt", sample_length, 0); + ReadHelper readHelper("normal_distribution.txt", sample_length, + population_length); readHelper.readFile(); std::vector sample, population; @@ -96,17 +97,18 @@ int main(int argc, char const *argv[]) { auto sort = CustomSort(sample_length, sizeof(key_type) * 8); - for (auto element : readHelper.population_vector) { + for (auto element : population) { auto ret = sort.sample_cdf_custom_version(sample.data(), element); auto cdf_index = ret / scale_size; auto index = trailing_zeroes((size_t)cdf_index); result[index]++; } - if (!multiple_times) + if (!multiple_times) { generateRandomLevel(result2, population_length); - else + } else { multiple_generate(result2, population_length); + } for (int i = 0; i < 32; i++) { if (!result[i] && !result2[i]) { -- cgit v1.3.1