From be0483ef4b6935324f85f124da8f4c1b93e918f6 Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Wed, 24 May 2023 14:12:56 +0800 Subject: fix(exp): Fix expt_0520 malloc failure Signed-off-by: KunoiSayami --- expt_0520.cu | 346 ++++++++++------------------------------------------------- 1 file changed, 56 insertions(+), 290 deletions(-) (limited to 'expt_0520.cu') diff --git a/expt_0520.cu b/expt_0520.cu index 839fef8..88815a7 100644 --- a/expt_0520.cu +++ b/expt_0520.cu @@ -64,12 +64,11 @@ 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 "read_helper.h" +#include "skiplistcustom.cuh" #include #include #include -#include #include #if __WORDSIZE == 64 @@ -116,11 +115,6 @@ 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 @@ -164,283 +158,116 @@ inline void __cudaCheckError(const char *file, const int line) { #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"); - } -} +__global__ void init(Node **n) { nodes = n; } // The main kernel -__global__ void kernel(LL *items, LL *op, LL *result) { +__global__ void kernel(LockFreeSkipList *skipList, const key_type *population, + size_t insertion_length) { // 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; + for (int 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) + auto tid = + i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x; + if (tid >= insertion_length) 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 - } + key_type item = population[tid]; + skipList->Add(item); } } -/*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., " + printf("Need two arguments: sample_length add insertion_length 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 + auto error = cudaGetLastError(); + auto total_row = 0UL; + auto sample_length = strtol(argv[1], nullptr, 0); + auto insertion_length = strtol(argv[2], nullptr, 0); + printf("insert_length %ld, sample_length: %ld\n", insertion_length, + sample_length); - // 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 - } + ReadHelper reader("normal_distribution.txt", sample_length, insertion_length); + printf("%d\n", __LINE__); + reader.readFile(total_row); - // 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; - } + printf("%d\n", __LINE__); + key_type *cudaPopulation; - adds = (NUM_ITEMS * adds) / 100; + printf("%d\n", __LINE__); + cudaMalloc(&cudaPopulation, sizeof(key_type) * insertion_length); + std::vector _sample, _population; + reader.split_into(_sample, _population); + cudaMemcpy(cudaPopulation, _population.data(), + sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice); + printf("%d\n", __LINE__); // 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; + Node **pointers = + (Node **)new LL[insertion_length]; // malloc(sizeof(LL) * adds); + Node **cudaNodePointers; + + printf("%d\n", __LINE__); // Allocate the pool of free nodes - for (i = 0; i < adds; i++) { - cudaMalloc((void **)&pointers[i], sizeof(Node)); + for (int i = 0; i < insertion_length; i++) { + cudaMalloc(&pointers[i], sizeof(Node)); } - cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds); - cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds, + cudaMalloc((void **)&cudaNodePointers, sizeof(Node *) * insertion_length); + cudaMemcpy(cudaNodePointers, pointers, sizeof(Node *) * insertion_length, cudaMemcpyHostToDevice); + CudaCheckError(); + printf("%d\n", __LINE__); // Allocate the skip list LockFreeSkipList *Clist; - auto *list = new LockFreeSkipList(levels); + auto *list = new LockFreeSkipList(_sample.data(), SAMPLE_SIZE); cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList)); cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice); + CudaCheckError(); // 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; + size_t blocks = (insertion_length % (NUM_THREADS * FACTOR) == 0) + ? insertion_length / (NUM_THREADS * FACTOR) + : (insertion_length / (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); + init<<<1, 32>>>(cudaNodePointers); cudaDeviceSynchronize(); // Launch main kernel @@ -450,13 +277,8 @@ int main(int argc, char **argv) { cudaEventCreate(&stop); cudaEventRecord(start, nullptr); - kernel<<>>(Citems, Cop, Cresult); + kernel<<>>(Clist, cudaPopulation, insertion_length); CudaCheckError(); - error = cudaGetLastError(); - if (cudaSuccess != error) { - printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error)); - // exit(-1); - } cudaDeviceSynchronize(); cudaEventRecord(stop, nullptr); cudaEventSynchronize(stop); @@ -471,58 +293,6 @@ int main(int argc, char **argv) { 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>>>(); @@ -552,16 +322,12 @@ int main(int argc, char **argv) { // 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;*/ + cudaFree(Clist); + for (int i = 0; i < insertion_length; i++) { + cudaFree(pointers[i]); + } + cudaFree(cudaNodePointers); + delete list; + delete[] pointers; return 0; } -- cgit v1.3.1