format: Run clang-format

Signed-off-by: KunoiSayami <[email protected]>

1 files changed, 377 insertions, 382 deletions

diff --git a/main.cu b/main.cu
index 74543ad..7c0aae3 100644
--- a/main.cu
+++ b/main.cu
@@ -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;