diff options
Diffstat (limited to 'main.cu')
| -rw-r--r-- | main.cu | 759 |
1 files changed, 377 insertions, 382 deletions
@@ -31,38 +31,43 @@ 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. + 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 + 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 + 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. + 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 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. + 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. + 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"stdio.h" -#include"stdlib.h" -#include"time.h" -#include"assert.h" +#include "cuda_runtime.h" +#include <cassert> +#include <cstdio> +#include <cstdlib> +#include <ctime> #if __WORDSIZE == 64 typedef unsigned long long LL; @@ -72,15 +77,15 @@ typedef unsigned int LL; // Maximum level of a node in the skip list //#define MAX_LEVEL 32 -constexpr size_t MAX_LEVEL=32; +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 = 512; -constexpr size_t NUM_ITEMS=1048576; -constexpr size_t KEYS=1048576; -constexpr size_t FACTOR=1; +constexpr size_t NUM_ITEMS = 1048576; +constexpr size_t KEYS = 1048576; +constexpr size_t FACTOR = 1; // Supported operations #define ADD (0) @@ -89,294 +94,276 @@ constexpr size_t FACTOR=1; #define CUDA_ERROR_CHECK -#define CudaSafeCall( err ) __cudaSafeCall( err, __FILE__, __LINE__ ) -#define CudaCheckError() __cudaCheckError( __FILE__, __LINE__ ) +#define CudaSafeCall(err) __cudaSafeCall(err, __FILE__, __LINE__) +#define CudaCheckError() __cudaCheckError(__FILE__, __LINE__) -inline void __cudaSafeCall( cudaError err, const char *file, const int 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 ); - } + 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 ) -{ +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 ); - } + 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 ); - } + // 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 __attribute__((aligned (16))) Node -{ - public: - int topLevel; // Level of the node - LL key; // Key value - LL next[MAX_LEVEL+1]; // Array of next links +class __attribute__((aligned(16))) 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 - __device__ __host__ LL CreateRef(Node* ref, bool mark) - { - LL val=(LL)ref; - val=val|mark; - return val; - } + // Create a next field from a reference and mark bit + __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); - } + __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 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); - } + // 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; - } + // 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 - Node(LL k) - { - key=k; - topLevel=MAX_LEVEL; - int i; - for(i=0;i<MAX_LEVEL+1;i++){ - next[i]=CreateRef((Node*)NULL, false); - } + // Constructor for sentinel nodes + 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 -{ - public: - Node* head; - Node* tail; - LockFreeSkipList() - { - Node* h=new Node(0); +class LockFreeSkipList { +public: + Node *head; + Node *tail; + LockFreeSkipList() { + Node *h = new Node(0); #if __WORDSIZE == 64 - Node* t=new Node((LL) NUM_ITEMS+10); + Node *t = new Node((LL)NUM_ITEMS + 10); #else - Node* t=new Node((LL)0xffffffff); + Node *t = new Node((LL)0xffffffff); #endif #ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void**)&head, sizeof(Node))); + CUDA_SAFE_CALL(cudaMalloc((void **)&head, sizeof(Node))); #else - cudaMalloc((void**)&head, sizeof(Node)); + cudaMalloc((void **)&head, sizeof(Node)); #endif #ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void**)&tail, sizeof(Node))); + CUDA_SAFE_CALL(cudaMalloc((void **)&tail, sizeof(Node))); #else - cudaMalloc((void**)&tail, sizeof(Node)); + cudaMalloc((void **)&tail, sizeof(Node)); #endif - int i; - for(i=0;i<h->topLevel+1;i++){ - h->SetRef(i, tail, false); - } + int i; + for (i = 0; i < h->topLevel + 1; i++) { + h->SetRef(i, tail, false); + } #ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice)); #else - cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice); + cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice); #endif #ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice)); #else - cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice); + cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice); #endif - } - __device__ bool find(LL, Node**, Node**); // Helping method - __device__ bool Add(LL); - __device__ bool Delete(LL); - __device__ bool Search(LL); + } + __device__ bool find(LL, Node **, Node **); // Helping method + __device__ bool Add(LL); + __device__ bool Delete(LL); + __device__ bool Search(LL); }; -__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__ 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 // Function for creating a new node when requested by an add operation -__device__ Node* GetNewNode(LL key) -{ - LL ind=atomicInc(&pointerIndex, NUM_ITEMS); - Node* n=nodes[ind]; - n->key=key; - n->topLevel=randoms[ind]; +__device__ Node *GetNewNode(LL key) { + LL ind = atomicInc(&pointerIndex, NUM_ITEMS); + Node *n = nodes[ind]; + n->key = key; + n->topLevel = randoms[ind]; int i; - for(i=0;i<n->topLevel+1;i++){ - n->SetRef(i, NULL, false); + for (i = 0; i < n->topLevel + 1; i++) { + n->SetRef(i, nullptr, false); } return n; } -__device__ LockFreeSkipList* l; // The lock-free skip list +__device__ LockFreeSkipList *l; // The lock-free skip list // Kernel for initializing device memory -__global__ void init(LockFreeSkipList* l1, Node** n, LL* rands) -{ - randoms=rands; - nodes=n; - l=l1; +__global__ void init(LockFreeSkipList *l1, Node **n, LL *rands) { + randoms = rands; + nodes = n; + l = l1; } // 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}; +LockFreeSkipList::find(LL key, Node **preds, + Node **succs) { // preds and succs are arrays of pointers + int bottomLevel = 0; + bool marked[] = {false}; bool snip = false; - Node* pred=NULL; - Node* curr=NULL; - Node* succ=NULL; - bool beenThereDoneThat= false; - while(true){ + Node *pred = nullptr; + Node *curr = nullptr; + Node *succ = nullptr; + bool beenThereDoneThat = false; + while (true) { beenThereDoneThat = false; - pred=head; + 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); + 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); + if (!snip) + break; + curr = pred->GetReference(level); + succ = curr->Get(level, marked); beenThereDoneThat = false; - //printf("find key is %d \n",(int)key); - } - if (beenThereDoneThat && !snip) break; - if(curr->key<=key){ - pred=curr; - curr=succ; + // printf("find key is %d \n",(int)key); } - else{ + if (beenThereDoneThat && !snip) + break; + if (curr->key <= key) { + pred = curr; + curr = succ; + } else { break; } } - if (beenThereDoneThat && !snip) break; - preds[level]=pred; - succs[level]=curr; + if (beenThereDoneThat && !snip) + break; + preds[level] = pred; + succs[level] = curr; } - if (beenThereDoneThat && !snip) continue; - return((curr->key==key)); + if (beenThereDoneThat && !snip) + continue; + return ((curr->key == key)); } } -__device__ bool -LockFreeSkipList::Search(LL key) -{ - int bottomLevel=0; - bool marked=false; - Node* pred=head; - Node* curr=NULL; - Node* succ=NULL; +__device__ bool LockFreeSkipList::Search(LL key) { + int bottomLevel = 0; + bool marked = false; + Node *pred = head; + Node *curr = nullptr; + Node *succ = nullptr; int level; - for(level=MAX_LEVEL;level>=bottomLevel;level--){ - curr=pred->GetReference(level); - while(true){ - succ=curr->Get(level, &marked); - while(marked){ - curr=curr->GetReference(level); - succ=curr->Get(level, &marked); - } - if(curr->key<key){ - pred=curr; - curr=succ; + for (level = MAX_LEVEL; level >= bottomLevel; level--) { + curr = pred->GetReference(level); + while (true) { + succ = curr->Get(level, &marked); + while (marked) { + curr = curr->GetReference(level); + succ = curr->Get(level, &marked); } - else{ + if (curr->key < key) { + pred = curr; + curr = succ; + } else { break; } } } - return(curr->key==key); + return (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){ +__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]; + } else { + Node *nodeToDelete = succs[bottomLevel]; int level; - for(level=nodeToDelete->topLevel;level>=bottomLevel+1;level--){ - succ=nodeToDelete->Get(level, marked); - while(marked[0]==false){ + 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(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==true){ + 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); return true; - } - else if(marked[0]==true){ + } else if (marked[0]) { return false; } } @@ -384,46 +371,46 @@ LockFreeSkipList::Delete(LL key) } } -__device__ bool -LockFreeSkipList::Add(LL key) -{ - Node* newNode=GetNewNode(key); - int topLevel=newNode->topLevel; - int bottomLevel=0; - Node* preds[MAX_LEVEL+1]; - Node* succs[MAX_LEVEL+1]; +__device__ bool LockFreeSkipList::Add(LL key) { + Node *newNode = GetNewNode(key); + 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){ + 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); + } 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]; + 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){ + // 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++){ + 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)){ + 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); + // 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); } } @@ -432,25 +419,25 @@ LockFreeSkipList::Add(LL key) } } -__global__ void print() -{ +__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!=NULL){ + 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); + 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); + 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); + 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"); } @@ -458,152 +445,159 @@ __global__ void print() // The main kernel -__global__ void kernel(LL* items, LL* op, LL* result) -{ +__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; + 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); + 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] == DELETE) { + result[tid] = l->Delete(item); } - if(op[tid]==SEARCH){ - result[tid]=l->Search(item); + if (op[tid] == SEARCH) { + result[tid] = l->Search(item); } } } // Generate the level of a newly created node -LL Randomlevel() -{ - LL v=1; - double p=0.5; - while(((rand()/(double)(RAND_MAX))<p) && (v<MAX_LEVEL)) v++; +LL Randomlevel() { + LL v = 1; + double p = 0.5; + while (((rand() / (double)(RAND_MAX)) < p) && (v < MAX_LEVEL)) + v++; return v; } -int main(int argc, char** argv) -{ +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); + 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 - int adds=atoi(argv[1]); - int deletes=atoi(argv[2]); + int adds = atoi(argv[1]); + int deletes = atoi(argv[2]); - if (adds+deletes > 100) { - printf("Sum of add and delete precentages exceeds 100.\nAborting...\n"); - exit(1); + if (adds + deletes > 100) { + printf("Sum of add and delete precentages exceeds 100.\nAborting...\n"); + exit(1); } // Allocate necessary arrays - LL* op=(LL*)malloc(sizeof(LL)*NUM_ITEMS); - LL* levels=(LL*)malloc(sizeof(LL)*NUM_ITEMS); - LL* items=(LL*)malloc(sizeof(LL)*NUM_ITEMS); - LL* result=(LL*)malloc(sizeof(LL)*NUM_ITEMS); + LL *op = (LL *)malloc(sizeof(LL) * NUM_ITEMS); + LL *levels = (LL *)malloc(sizeof(LL) * NUM_ITEMS); + LL *items = (LL *)malloc(sizeof(LL) * NUM_ITEMS); + LL *result = (LL *)malloc(sizeof(LL) * NUM_ITEMS); int i; // NUM_ITEMS is the total number of operations to execute srand(0); - for(i=0;i<NUM_ITEMS;i++){ - items[i]=i+3;//10+rand()%KEYS; // Keys associated with operations + for (i = 0; i < NUM_ITEMS; i++) { + items[i] = i + 3; // 10+rand()%KEYS; // Keys associated with + // operations } - - for(i=0;i<NUM_ITEMS;i++){ - int first = rand()%NUM_ITEMS; - int second = rand()%NUM_ITEMS; + for (i = 0; i < NUM_ITEMS; i++) { + int first = rand() % NUM_ITEMS; + int second = rand() % NUM_ITEMS; LL temp; temp = items[first]; - items[first]=items[second]; - items[second]=temp; + items[first] = items[second]; + items[second] = temp; } // Pre-generated levels of skip list nodes (relevant only if op[i] is add) srand(0); - for(i=0;i<NUM_ITEMS;i++){ - levels[i]=Randomlevel()-1; + for (i = 0; i < NUM_ITEMS; i++) { + levels[i] = Randomlevel() - 1; } // Populate the sequence of operations - for(i=0;i<(NUM_ITEMS*adds)/100;i++){ - op[i]=ADD; + 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 * (adds + deletes)) / 100; i++) { + op[i] = DELETE; } - for(;i<NUM_ITEMS;i++){ - op[i]=SEARCH; + for (; i < NUM_ITEMS; i++) { + op[i] = SEARCH; } - adds=(NUM_ITEMS*adds)/100; + adds = (NUM_ITEMS * adds) / 100; // Allocate device memory - LL* Citems; - LL* Cop; - LL* Cresult; - LL* Clevels; + LL *Citems; + LL *Cop; + LL *Cresult; + LL *Clevels; #ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void**)&Cresult, sizeof(LL)*NUM_ITEMS)); - CUDA_SAFE_CALL(cudaMalloc((void**)&Citems, sizeof(LL)*NUM_ITEMS)); - CUDA_SAFE_CALL(cudaMalloc((void**)&Cop, sizeof(LL)*NUM_ITEMS)); - CUDA_SAFE_CALL(cudaMalloc((void**)&Clevels, sizeof(LL)*NUM_ITEMS)); - CUDA_SAFE_CALL(cudaMemcpy(Clevels, levels, sizeof(LL)*NUM_ITEMS, cudaMemcpyHostToDevice)); - CUDA_SAFE_CALL(cudaMemcpy(Citems, items, sizeof(LL)*NUM_ITEMS, cudaMemcpyHostToDevice)); - CUDA_SAFE_CALL(cudaMemcpy(Cop, op, sizeof(LL)*NUM_ITEMS, cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMalloc((void **)&Cresult, sizeof(LL) * NUM_ITEMS)); + CUDA_SAFE_CALL(cudaMalloc((void **)&Citems, sizeof(LL) * NUM_ITEMS)); + CUDA_SAFE_CALL(cudaMalloc((void **)&Cop, sizeof(LL) * NUM_ITEMS)); + CUDA_SAFE_CALL(cudaMalloc((void **)&Clevels, sizeof(LL) * NUM_ITEMS)); + CUDA_SAFE_CALL(cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS, + cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMemcpy(Citems, items, sizeof(LL) * NUM_ITEMS, + cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL( + cudaMemcpy(Cop, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice)); #else - 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); + 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); #endif - Node** pointers=(Node**)malloc(sizeof(LL)*adds); - Node** Cpointers; + Node **pointers = (Node **)malloc(sizeof(LL) * adds); + Node **Cpointers; // Allocate the pool of free nodes - for(i=0;i<adds;i++){ + for (i = 0; i < adds; i++) { #ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void**)&pointers[i], sizeof(Node))); + CUDA_SAFE_CALL(cudaMalloc((void **)&pointers[i], sizeof(Node))); #else - cudaMalloc((void**)&pointers[i], sizeof(Node)); + cudaMalloc((void **)&pointers[i], sizeof(Node)); #endif } #ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void**)&Cpointers, sizeof(Node*)*adds)); - CUDA_SAFE_CALL(cudaMemcpy(Cpointers,pointers, sizeof(Node*)*adds, cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds)); + CUDA_SAFE_CALL(cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds, + cudaMemcpyHostToDevice)); #else - cudaMalloc((void**)&Cpointers, sizeof(Node*)*adds); - cudaMemcpy(Cpointers,pointers, sizeof(Node*)*adds, cudaMemcpyHostToDevice); + cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds); + cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds, + cudaMemcpyHostToDevice); #endif // Allocate the skip list - LockFreeSkipList* Clist; - LockFreeSkipList* list=new LockFreeSkipList(); + LockFreeSkipList *Clist; + auto *list = new LockFreeSkipList(); #ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void**)&Clist, sizeof(LockFreeSkipList))); - CUDA_SAFE_CALL(cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList))); + CUDA_SAFE_CALL(cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), + cudaMemcpyHostToDevice)); #else - cudaMalloc((void**)&Clist, sizeof(LockFreeSkipList)); + cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList)); cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice); #endif @@ -612,35 +606,37 @@ int main(int argc, char** argv) // 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; + 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)); + 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, Clevels); + init<<<1, 32>>>(Clist, Cpointers, Clevels); cudaDeviceSynchronize(); // Launch main kernel - cudaEvent_t start,stop; + cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); - cudaEventRecord(start,0); + cudaEventRecord(start, 0); - kernel<<<blocks,NUM_THREADS>>>(Citems, Cop, Cresult); + kernel<<<blocks, NUM_THREADS>>>(Citems, Cop, Cresult); CudaCheckError(); - error=cudaGetLastError(); - if(cudaSuccess!=error){ - printf("error0:CUDA ERROR (%d) {%s}\n",error,cudaGetErrorString(error)); - //exit(-1); + error = cudaGetLastError(); + if (cudaSuccess != error) { + printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error)); + // exit(-1); } cudaDeviceSynchronize(); - cudaEventRecord(stop,0); + cudaEventRecord(stop, 0); cudaEventSynchronize(stop); float time; cudaEventElapsedTime(&time, start, stop); @@ -649,34 +645,32 @@ int main(int argc, char** argv) // Print kernel execution time in milliseconds - printf("%lf\n",time); + printf("%lf\n", time); // Launch main kernel for query // Populate the sequence of operations - for(i=0;i<NUM_ITEMS;i++){ - op[i]=SEARCH; + 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); - - + LL *Cop2; + cudaMalloc((void **)&Cop2, sizeof(LL) * NUM_ITEMS); + cudaMemcpy(Cop2, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice); cudaEventCreate(&start); cudaEventCreate(&stop); - cudaEventRecord(start,0); + cudaEventRecord(start, 0); - kernel<<<blocks,NUM_THREADS>>>(Citems, Cop2, Cresult); + kernel<<<blocks, NUM_THREADS>>>(Citems, Cop2, Cresult); CudaCheckError(); - error=cudaGetLastError(); - if(cudaSuccess!=error){ - printf("error0:CUDA ERROR (%d) {%s}\n",error,cudaGetErrorString(error)); - //exit(-1); + error = cudaGetLastError(); + if (cudaSuccess != error) { + printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error)); + // exit(-1); } cudaDeviceSynchronize(); - cudaEventRecord(stop,0); + cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&time, start, stop); cudaEventDestroy(start); @@ -684,26 +678,27 @@ int main(int argc, char** argv) // Print kernel execution time in milliseconds - printf("%lf\n",time); + printf("%lf\n", time); // Check for errors - error=cudaGetLastError(); - if(cudaSuccess!=error){ - printf("error1:CUDA ERROR (%d) {%s}\n",error, cudaGetErrorString(error)); + error = cudaGetLastError(); + if (cudaSuccess != error) { + printf("error1:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error)); exit(-1); } // Move results back to host memory #ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMemcpy(result, Cresult, sizeof(LL)*NUM_ITEMS, cudaMemcpyDeviceToHost)); + CUDA_SAFE_CALL(cudaMemcpy(result, Cresult, sizeof(LL) * NUM_ITEMS, + cudaMemcpyDeviceToHost)); #else - cudaMemcpy(result, Cresult, sizeof(LL)*NUM_ITEMS, cudaMemcpyDeviceToHost); + cudaMemcpy(result, Cresult, sizeof(LL) * NUM_ITEMS, cudaMemcpyDeviceToHost); #endif // Uncomment the following for debugging - //print<<<1,32>>>(); + // print<<<1,32>>>(); cudaDeviceSynchronize(); return 0; |