From 70af7c287739bc45089af9c36e6a41a2a317d61a Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Mon, 22 May 2023 19:42:05 +0800 Subject: feat: Add skiplist header Signed-off-by: KunoiSayami --- expt_0520.cu | 567 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 567 insertions(+) create mode 100644 expt_0520.cu (limited to 'expt_0520.cu') diff --git a/expt_0520.cu b/expt_0520.cu new file mode 100644 index 0000000..839fef8 --- /dev/null +++ b/expt_0520.cu @@ -0,0 +1,567 @@ +/* + +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 "skiplist.cuh" +#include +#include +#include +#include +#include +#include + +#if __WORDSIZE == 64 +typedef unsigned long long LL; +#else +typedef unsigned int LL; +#endif + +#ifndef BUILD_SIZE +#define BUILD_SIZE 1048576 +#endif + +#ifndef STEP_SIZE +#define STEP_SIZE 2 +#endif + +// #define MEASURE_TIME +// #define MEASURE_ACCESS + +#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 = 16; + +// Number of threads per block +// #define NUM_THREADS 512 +constexpr size_t NUM_THREADS = 512; + +// 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 = 1024; + +constexpr int block_size = STEP_SIZE; + +// Supported operations +constexpr int ADD = 0; +constexpr int DELETE = 1; +constexpr int SEARCH = 2; + +typedef LL key_type; + +#ifdef RANDOM_TARGET +constexpr const char *TARGET_STRING = "RANDOM"; +#else +constexpr const char *TARGET_STRING = "PERFECT"; +#endif + +#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 +} + +__device__ LockFreeSkipList *l; // The lock-free skip list + +__device__ key_type SampleStorage[SAMPLE_SIZE]; + +// Kernel for initializing device memory + +__global__ void init(LockFreeSkipList *lockFreeSkipList, Node **n) { + nodes = n; + l = lockFreeSkipList; +} + +__global__ void print() { + // For debugging + int tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid == 0) { + Node *p = l->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(LL *items, LL *op, 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 + + int tid, i; + for (i = 0; i < FACTOR; + i++) { // FACTOR is the number of operations per thread + tid = i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x; + if (tid >= NUM_ITEMS) + return; + + // Grab the operation and the associated key and execute + LL item = items[tid]; + if (op[tid] == ADD) { + result[tid] = l->Add(item); + } + if (op[tid] == DELETE) { + result[tid] = l->Delete(item); + } + if (op[tid] == SEARCH) { +#ifdef MEASURE_TIME + unsigned long long start_time = clock64(); +#endif + result[tid] = l->Search(item); +#ifdef MEASURE_TIME + unsigned long long end_time = clock64() - start_time; + if (l->spend_time[tid]) { + printf("conflict: %d\n", tid); + } + l->spend_time[tid] = (int)end_time; +#endif + } + } +} + +/*LL Randomlevel() { + LL v = 1; + double p = 0.5; + while (((rand() / (double)(RAND_MAX)) < p) && (v < MAX_LEVEL)) + v++; + return v; +}*/ + +// Generate the level of a newly created node +LL RandomLevel(std::mt19937 &randomEngine, double p) { + std::geometric_distribution<> distribution(p); + return std::min(MAX_LEVEL, (size_t)distribution(randomEngine)); +} + +std::vector storage; + +unsigned trailing_zeroes(size_t index) { + unsigned bits = 0; + LL x = index / block_size; + + if (x) { + while (x % block_size == 0) { + ++bits; + x /= block_size; + } + } + return bits; +} + +LL CustomLevel(LL value) { + auto left = std::lower_bound(storage.begin(), storage.end(), value); + auto right = std::upper_bound(storage.begin(), storage.end(), value); + + if (right - left != 1) { + printf("%ld\n", right - left); + } + assert(right - left == 1); + + auto index = left - storage.begin(); + + if (index % block_size == 0) { + auto level = trailing_zeroes(index) + 1; + // printf("%ld,%u\n", index, level); + return level; + } + + return 1; +} + +__global__ void print_function() { +#ifdef MEASURE_ACCESS + printf("count: %u\n", l->getAccessCount()); +#endif +} + +std::vector population; +std::vector sample; + +void initialize(const std::vector &input_population, + const std::vector &input_sample) { + population = input_population; + std::sort(population.begin(), population.end()); + sample = input_sample; + std::sort(sample.begin(), sample.end()); +} + +int main(int argc, char **argv) { + if (argc != 3) { + printf("Need two arguments: percent add ops and percent delete ops (e.g., " + "30 50 for 30%% add and 50%% delete).\nAborting...\n"); + exit(1); + } + + // Extract operations ratio + long adds = strtol(argv[1], nullptr, 10); + long deletes = strtol(argv[2], nullptr, 10); + + storage.reserve(NUM_ITEMS); + + if (adds + deletes > 100) { + printf("Sum of add and delete percentages exceeds 100.\nAborting...\n"); + exit(1); + } + + // 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[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); + int i; + + // NUM_ITEMS is the total number of operations to execute + // srand(0); + + std::random_device randomDevice; + std::mt19937 randomEngine(randomDevice()); + std::uniform_int_distribution uniformIntDistributionArray(0, + NUM_ITEMS - 1); + // std::vector storage; + + // std::normal_distribution normalDistribution{2147483647, + // 2147483647}; + + for (i = 0; i < NUM_ITEMS; i++) { + items[i] = i + 3; // 10+rand()%KEYS; + // Keys associated with + // operations + storage.push_back(i + 3); + // auto key = (key_type)std::round(normalDistribution(randomEngine)); + // items[i] = key; + // storage.push_back(key); + } + + std::sort(storage.begin(), storage.end()); + +#if 0 + for (i = 0; i < NUM_ITEMS; i++) { + /*int first = rand() % NUM_ITEMS; + int second = rand() % NUM_ITEMS;*/ + + std::swap(items[uniformIntDistributionArray(randomEngine)], + items[uniformIntDistributionArray(randomEngine)]); + /*LL temp; + temp = items[first]; + items[first] = items[second]; + items[second] = temp;*/ + } +#endif + + // Pre-generated levels of skip list nodes (relevant only if op[i] is add) + // srand(0); + for (i = 0; i < NUM_ITEMS; i++) { +#ifdef RANDOM_HEIGHT + levels[i] = RandomLevel(1 / randomEngine) - 1; // 36/14 +#else + levels[i] = CustomLevel(items[i]) - 1; // 31/18 +#endif + } + + // Populate the sequence of operations + for (i = 0; i < (NUM_ITEMS * adds) / 100; i++) { + op[i] = ADD; + } + for (; i < (NUM_ITEMS * (adds + deletes)) / 100; i++) { + op[i] = DELETE; + } + for (; i < NUM_ITEMS; i++) { + op[i] = SEARCH; + } + + adds = (NUM_ITEMS * adds) / 100; + + // Allocate device memory + + LL *Citems; + LL *Cop; + LL *Cresult; + // LL *Clevels; + + cudaMalloc((void **)&Cresult, sizeof(LL) * NUM_ITEMS); + cudaMalloc((void **)&Citems, sizeof(LL) * NUM_ITEMS); + cudaMalloc((void **)&Cop, sizeof(LL) * NUM_ITEMS); + // cudaMalloc((void **)&Clevels, sizeof(LL) * NUM_ITEMS); + // cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS, + // cudaMemcpyHostToDevice); + cudaMemcpy(Citems, items, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice); + cudaMemcpy(Cop, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice); + Node **pointers = (Node **)new LL[adds]; // malloc(sizeof(LL) * adds); + Node **Cpointers; + + // Allocate the pool of free nodes + + for (i = 0; i < adds; i++) { + cudaMalloc((void **)&pointers[i], sizeof(Node)); + } + cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds); + cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds, + cudaMemcpyHostToDevice); + + // Allocate the skip list + + LockFreeSkipList *Clist; + auto *list = new LockFreeSkipList(levels); + + cudaMalloc((void **)&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 + + int blocks = (NUM_ITEMS % (NUM_THREADS * FACTOR) == 0) + ? NUM_ITEMS / (NUM_THREADS * FACTOR) + : (NUM_ITEMS / (NUM_THREADS * FACTOR)) + 1; + + // Error checking code + cudaError_t error = cudaGetLastError(); + if (cudaSuccess != error) { + printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error)); + exit(-1); + } + + // Initialize the device memory + init<<<1, 32>>>(Clist, Cpointers); + cudaDeviceSynchronize(); + + // Launch main kernel + + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start, nullptr); + + kernel<<>>(Citems, Cop, Cresult); + CudaCheckError(); + error = cudaGetLastError(); + if (cudaSuccess != error) { + printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error)); + // exit(-1); + } + cudaDeviceSynchronize(); + cudaEventRecord(stop, nullptr); + cudaEventSynchronize(stop); + float time; + cudaEventElapsedTime(&time, start, stop); + cudaEventDestroy(start); + cudaEventDestroy(stop); + + // Print kernel execution time in milliseconds + + printf("%s %d ", TARGET_STRING, block_size); + + printf("%lu: %lf", NUM_ITEMS, time); + + // Launch main kernel for query + + // Populate the sequence of operations + for (i = 0; i < NUM_ITEMS; i++) { + op[i] = SEARCH; + } + + LL *Cop2; + cudaMalloc((void **)&Cop2, sizeof(LL) * NUM_ITEMS); + cudaMemcpy(Cop2, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice); + + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start, nullptr); +#ifdef MEASURE_TIME + kernel<<>>(Citems, Cop2, Cresult); +#else + kernel<<>>(Citems, Cop2, Cresult); +#endif + CudaCheckError(); + error = cudaGetLastError(); + if (cudaSuccess != error) { + printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error)); + // exit(-1); + } + cudaDeviceSynchronize(); + cudaEventRecord(stop, nullptr); + cudaEventSynchronize(stop); + cudaEventElapsedTime(&time, start, stop); + cudaEventDestroy(start); + cudaEventDestroy(stop); + + // Print kernel execution time in milliseconds + + printf(" %lf\n", time); + + // Check for errors + + error = cudaGetLastError(); + if (cudaSuccess != error) { + printf("error1:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error)); + exit(-1); + } + + // Move results back to host memory + + cudaMemcpy(result, Cresult, sizeof(LL) * NUM_ITEMS, cudaMemcpyDeviceToHost); + + // Uncomment the following for debugging + // print<<<1,32>>>(); + 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 (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(Cop2); + cudaFree(Clevels); + cudaFree(Cop); + cudaFree(Citems); + cudaFree(Cresult); + free(pointers); + delete [] op; + delete [] levels; + delete [] items; + delete [] result;*/ + return 0; +} -- cgit v1.3.1