From 9de90473c345d7ce5b4381104ae6807b1ebad36f Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Fri, 27 May 2022 23:06:02 +0800 Subject: format: Run clang-format Signed-off-by: KunoiSayami --- main.cu | 763 ++++++++++++++++++++++++++++++++-------------------------------- 1 file changed, 379 insertions(+), 384 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 +#include +#include +#include #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 - - // 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; - } +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; + } - __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;itopLevel+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;itopLevel+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= 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++){ - - 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)){ + for (level = bottomLevel + 1; level <= topLevel; level++) { + + while (true) { + pred = preds[level]; + succ = succs[level]; + newNode->SetRef(level, succ, false); + // printf("-- key is %d pred is %d succ is %d level is %d + // \n",(int)key,(int)pred->key,(int)succ->key,(int)level); + if (pred->CompareAndSet(level, succ, newNode, false, false)) { break; } - //printf("key is %d pred is %d succ is %d level is %d \n",(int)key,(int)pred->key,(int)succ->key,(int)level); + // 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=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)) 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>>(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<<>>(Citems, Cop, Cresult); + kernel<<>>(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>>(Citems, Cop2, Cresult); + kernel<<>>(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; -- cgit v1.3.1