diff options
Diffstat (limited to 'benchmarks/CUDA/NQU/nqueen.cu')
| -rw-r--r-- | benchmarks/CUDA/NQU/nqueen.cu | 758 |
1 files changed, 758 insertions, 0 deletions
diff --git a/benchmarks/CUDA/NQU/nqueen.cu b/benchmarks/CUDA/NQU/nqueen.cu new file mode 100644 index 0000000..68dc9b5 --- /dev/null +++ b/benchmarks/CUDA/NQU/nqueen.cu @@ -0,0 +1,758 @@ +// N-queen for CUDA
+//
+// Copyright(c) 2008 Ping-Che Chen
+
+
+//#define WIN32_LEAN_AND_MEAN
+//#include <windows.h>
+#include <stdio.h>
+#include <cutil.h>
+
+#define THREAD_NUM 96
+
+
+int bunk = 0; // this is a dummy variable used for making sure clock() are not optimized out
+
+/*
+ * ----------------------------------------------------------------
+ * This is a recursive version of n-queen backtracking solver.
+ * A non-recursive version is used instead.
+ * ----------------------------------------------------------------
+
+long long solve_nqueen_internal(int n, unsigned int mask, unsigned int l_mask, unsigned int r_mask, unsigned int t_mask)
+{
+ if(mask == t_mask) {
+ return 1;
+ }
+
+ unsigned int m = (mask | l_mask | r_mask);
+ if((m & t_mask) == t_mask) {
+ return 0;
+ }
+
+ long long total = 0;
+ unsigned int index = (m + 1) & ~m;
+ while((index & t_mask) != 0) {
+ total += solve_nqueen_internal(mask | index, (l_mask | index) << 1, (r_mask | index) >> 1, t_mask);
+ m |= index;
+ index = (m + 1) & ~m;
+ }
+
+ return total;
+}
+
+
+long long solve_nqueen(int n)
+{
+ return solve_nqueen_internal(0, 0, 0, (1 << n) - 1);
+}
+*/
+
+/* -------------------------------------------------------------------
+ * This is a non-recursive version of n-queen backtracking solver.
+ * This provides the basis for the CUDA version.
+ * -------------------------------------------------------------------
+ */
+
+long long solve_nqueen(int n)
+{
+ unsigned int mask[32];
+ unsigned int l_mask[32];
+ unsigned int r_mask[32];
+ unsigned int m[32];
+
+ if(n <= 0 || n > 32) {
+ return 0;
+ }
+
+ const unsigned int t_mask = (1 << n) - 1;
+ long long total = 0;
+ long long upper_total = 0;
+ int i = 0, j;
+ unsigned int index;
+
+ mask[0] = 0;
+ l_mask[0] = 0;
+ r_mask[0] = 0;
+ m[0] = 0;
+
+ for(j = 0; j < (n + 1) / 2; j++) {
+ index = (1 << j);
+ m[0] |= index;
+
+ mask[1] = index;
+ l_mask[1] = index << 1;
+ r_mask[1] = index >> 1;
+ m[1] = (mask[1] | l_mask[1] | r_mask[1]);
+ i = 1;
+
+ if(n % 2 == 1 && j == (n + 1) / 2 - 1) {
+ upper_total = total;
+ total = 0;
+ }
+
+ while(i > 0) {
+ if((m[i] & t_mask) == t_mask) {
+ i--;
+ }
+ else {
+ index = ((m[i] + 1) ^ m[i]) & ~m[i];
+ m[i] |= index;
+ if((index & t_mask) != 0) {
+ if(i + 1 == n) {
+ total++;
+ i--;
+ }
+ else {
+ mask[i + 1] = mask[i] | index;
+ l_mask[i + 1] = (l_mask[i] | index) << 1;
+ r_mask[i + 1] = (r_mask[i] | index) >> 1;
+ m[i + 1] = (mask[i + 1] | l_mask[i + 1] | r_mask[i + 1]);
+ i++;
+ }
+ }
+ else {
+ i --;
+ }
+ }
+ }
+ }
+
+ bunk = 2;
+
+ if(n % 2 == 0) {
+ return total * 2;
+ }
+ else {
+ return upper_total * 2 + total;
+ }
+}
+
+
+/* -------------------------------------------------------------------
+ * This is a non-recursive version of n-queen backtracking solver
+ * with multi-thread support.
+ * -------------------------------------------------------------------
+ */
+/*
+struct thread_context
+{
+ HANDLE thread;
+ bool stop;
+
+ long long total;
+ int n;
+ unsigned int mask;
+ unsigned int l_mask;
+ unsigned int r_mask;
+ unsigned int t_mask;
+
+ HANDLE ready;
+ HANDLE complete;
+};
+
+DWORD WINAPI solve_nqueen_proc(LPVOID param)
+{
+ thread_context* ctx = (thread_context*) param;
+
+ unsigned int mask[32];
+ unsigned int l_mask[32];
+ unsigned int r_mask[32];
+ unsigned int m[32];
+ unsigned int t_mask;
+ long long total;
+ unsigned int index;
+ unsigned int mark;
+
+ for(;;) {
+ WaitForSingleObject(ctx->ready, INFINITE);
+ if(ctx->stop) {
+ break;
+ }
+
+ int i = 0;
+
+ mask[0] = ctx->mask;
+ l_mask[0] = ctx->l_mask;
+ r_mask[0] = ctx->r_mask;
+ m[0] = mask[0] | l_mask[0] | r_mask[0];
+ total = 0;
+ t_mask = ctx->t_mask;
+ mark = ctx->n;
+
+ while(i >= 0) {
+ if((m[i] & t_mask) == t_mask) {
+ i--;
+ }
+ else {
+ index = (m[i] + 1) & ~m[i];
+ m[i] |= index;
+ if((index & t_mask) != 0) {
+ if(i + 1 == mark) {
+ total++;
+ i--;
+ }
+ else {
+ mask[i + 1] = mask[i] | index;
+ l_mask[i + 1] = (l_mask[i] | index) << 1;
+ r_mask[i + 1] = (r_mask[i] | index) >> 1;
+ m[i + 1] = (mask[i + 1] | l_mask[i + 1] | r_mask[i + 1]);
+ i++;
+ }
+ }
+ else {
+ i --;
+ }
+ }
+ }
+
+ ctx->total = total;
+
+ SetEvent(ctx->complete);
+ }
+
+ return 0;
+}
+
+long long solve_nqueen_mcpu(int n)
+{
+ if(n <= 0 || n > 32) {
+ return 0;
+ }
+
+ SYSTEM_INFO info;
+ thread_context* threads;
+ int num_threads;
+
+ GetSystemInfo(&info);
+ num_threads = info.dwNumberOfProcessors;
+ if(num_threads == 1) {
+ // only one cpu found, use single thread version
+ return solve_nqueen(n);
+ }
+
+ threads = new thread_context[num_threads];
+ int j;
+ for(j = 0; j < num_threads; j++) {
+ threads[j].stop = false;
+ threads[j].ready = CreateEvent(0, FALSE, FALSE, 0);
+ threads[j].complete = CreateEvent(0, FALSE, TRUE, 0);
+ threads[j].thread = CreateThread(0, 0, solve_nqueen_proc, threads + j, 0, 0);
+ threads[j].total = 0;
+ }
+
+ int thread_idx = 0;
+
+ const unsigned int t_mask = (1 << n) - 1;
+ long long total = 0;
+ unsigned int index;
+
+ unsigned int m_mask = 0;
+ if(n % 2 == 1) {
+ m_mask = 1 << ((n + 1) / 2 - 1);
+ }
+
+ for(j = 0; j < (n + 1) / 2; j++) {
+ index = 1 << j;
+
+ WaitForSingleObject(threads[thread_idx].complete, INFINITE);
+
+ if(threads[thread_idx].mask != m_mask) {
+ total += threads[thread_idx].total * 2;
+ }
+ else {
+ total += threads[thread_idx].total;
+ }
+
+ threads[thread_idx].mask = index;
+ threads[thread_idx].l_mask = index << 1;
+ threads[thread_idx].r_mask = index >> 1;
+ threads[thread_idx].t_mask = t_mask;
+ threads[thread_idx].total = 0;
+ threads[thread_idx].n = n - 1;
+
+ SetEvent(threads[thread_idx].ready);
+
+ thread_idx = (thread_idx + 1) % num_threads;
+ }
+
+ // collect all threads...
+ HANDLE* events = new HANDLE[num_threads];
+ for(j = 0; j < num_threads; j++) {
+ events[j] = threads[j].complete;
+ }
+ WaitForMultipleObjects(num_threads, events, TRUE, INFINITE);
+ for(j = 0; j < num_threads; j++) {
+ if(threads[j].mask != m_mask) {
+ total += threads[j].total * 2;
+ }
+ else {
+ total += threads[j].total;
+ }
+
+ threads[j].stop = true;
+ SetEvent(threads[j].ready);
+
+ events[j] = threads[j].thread;
+ }
+
+ WaitForMultipleObjects(num_threads, events, TRUE, INFINITE);
+
+ for(j = 0; j < num_threads; j++) {
+ CloseHandle(threads[j].thread);
+ CloseHandle(threads[j].ready);
+ CloseHandle(threads[j].complete);
+ }
+ delete[] threads;
+ delete[] events;
+
+ bunk = 3;
+
+ return total;
+}
+*/
+
+
+/* --------------------------------------------------------------------------
+ * This is a non-recursive version of n-queen backtracking solver for CUDA.
+ * It receives multiple initial conditions from a CPU iterator, and count
+ * each conditions.
+ * --------------------------------------------------------------------------
+ */
+
+__global__ void solve_nqueen_cuda_kernel(int n, int mark, unsigned int* total_masks, unsigned int* total_l_masks, unsigned int* total_r_masks, unsigned int* results, int total_conditions)
+{
+ const int tid = threadIdx.x;
+ const int bid = blockIdx.x;
+ const int idx = bid * blockDim.x + tid;
+
+ __shared__ unsigned int mask[THREAD_NUM][10];
+ __shared__ unsigned int l_mask[THREAD_NUM][10];
+ __shared__ unsigned int r_mask[THREAD_NUM][10];
+ __shared__ unsigned int m[THREAD_NUM][10];
+
+ __shared__ unsigned int sum[THREAD_NUM];
+
+ const unsigned int t_mask = (1 << n) - 1;
+ int total = 0;
+ int i = 0;
+ unsigned int index;
+
+ if(idx < total_conditions) {
+ mask[tid][i] = total_masks[idx];
+ l_mask[tid][i] = total_l_masks[idx];
+ r_mask[tid][i] = total_r_masks[idx];
+ m[tid][i] = mask[tid][i] | l_mask[tid][i] | r_mask[tid][i];
+
+ while(i >= 0) {
+ if((m[tid][i] & t_mask) == t_mask) {
+ i--;
+ }
+ else {
+ index = (m[tid][i] + 1) & ~m[tid][i];
+ m[tid][i] |= index;
+ if((index & t_mask) != 0) {
+ if(i + 1 == mark) {
+ total++;
+ i--;
+ }
+ else {
+ mask[tid][i + 1] = mask[tid][i] | index;
+ l_mask[tid][i + 1] = (l_mask[tid][i] | index) << 1;
+ r_mask[tid][i + 1] = (r_mask[tid][i] | index) >> 1;
+ m[tid][i + 1] = (mask[tid][i + 1] | l_mask[tid][i + 1] | r_mask[tid][i + 1]);
+ i++;
+ }
+ }
+ else {
+ i --;
+ }
+ }
+ }
+
+ sum[tid] = total;
+ }
+ else {
+ sum[tid] = 0;
+ }
+
+ __syncthreads();
+
+ // reduction
+ if(tid < 64 && tid + 64 < THREAD_NUM) { sum[tid] += sum[tid + 64]; } __syncthreads();
+ if(tid < 32) { sum[tid] += sum[tid + 32]; } __syncthreads();
+ if(tid < 16) { sum[tid] += sum[tid + 16]; } __syncthreads();
+ if(tid < 8) { sum[tid] += sum[tid + 8]; } __syncthreads();
+ if(tid < 4) { sum[tid] += sum[tid + 4]; } __syncthreads();
+ if(tid < 2) { sum[tid] += sum[tid + 2]; } __syncthreads();
+ if(tid < 1) { sum[tid] += sum[tid + 1]; } __syncthreads();
+
+ if(tid == 0) {
+ results[bid] = sum[0];
+ }
+}
+
+
+long long solve_nqueen_cuda(int n, int steps)
+{
+ // generating start conditions
+ unsigned int mask[32];
+ unsigned int l_mask[32];
+ unsigned int r_mask[32];
+ unsigned int m[32];
+ unsigned int index;
+
+ if(n <= 0 || n > 32) {
+ return 0;
+ }
+
+ unsigned int* total_masks = new unsigned int[steps];
+ unsigned int* total_l_masks = new unsigned int[steps];
+ unsigned int* total_r_masks = new unsigned int[steps];
+ unsigned int* results = new unsigned int[steps];
+
+ unsigned int* masks_cuda;
+ unsigned int* l_masks_cuda;
+ unsigned int* r_masks_cuda;
+ unsigned int* results_cuda;
+
+ cudaMalloc((void**) &masks_cuda, sizeof(int) * steps);
+ cudaMalloc((void**) &l_masks_cuda, sizeof(int) * steps);
+ cudaMalloc((void**) &r_masks_cuda, sizeof(int) * steps);
+ cudaMalloc((void**) &results_cuda, sizeof(int) * steps / THREAD_NUM);
+
+ const unsigned int t_mask = (1 << n) - 1;
+ const unsigned int mark = n > 11 ? n - 10 : 2;
+ long long total = 0;
+ int total_conditions = 0;
+ int i = 0, j;
+
+ mask[0] = 0;
+ l_mask[0] = 0;
+ r_mask[0] = 0;
+ m[0] = 0;
+
+ bool computed = false;
+
+ for(j = 0; j < n / 2; j++) {
+ index = (1 << j);
+ m[0] |= index;
+
+ mask[1] = index;
+ l_mask[1] = index << 1;
+ r_mask[1] = index >> 1;
+ m[1] = (mask[1] | l_mask[1] | r_mask[1]);
+ i = 1;
+
+ while(i > 0) {
+ if((m[i] & t_mask) == t_mask) {
+ i--;
+ }
+ else {
+ index = (m[i] + 1) & ~m[i];
+ m[i] |= index;
+ if((index & t_mask) != 0) {
+ mask[i + 1] = mask[i] | index;
+ l_mask[i + 1] = (l_mask[i] | index) << 1;
+ r_mask[i + 1] = (r_mask[i] | index) >> 1;
+ m[i + 1] = (mask[i + 1] | l_mask[i + 1] | r_mask[i + 1]);
+ i++;
+ if(i == mark) {
+ total_masks[total_conditions] = mask[i];
+ total_l_masks[total_conditions] = l_mask[i];
+ total_r_masks[total_conditions] = r_mask[i];
+ total_conditions++;
+ if(total_conditions == steps) {
+ if(computed) {
+ cudaMemcpy(results, results_cuda, sizeof(int) * steps / THREAD_NUM, cudaMemcpyDeviceToHost);
+
+ for(int j = 0; j < steps / THREAD_NUM; j++) {
+ total += results[j];
+ }
+
+ computed = false;
+ }
+
+ // start computation
+ cudaMemcpy(masks_cuda, total_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+ cudaMemcpy(l_masks_cuda, total_l_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+ cudaMemcpy(r_masks_cuda, total_r_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+
+ solve_nqueen_cuda_kernel<<<steps/THREAD_NUM, THREAD_NUM>>>(n, n - mark, masks_cuda, l_masks_cuda, r_masks_cuda, results_cuda, total_conditions);
+
+ computed = true;
+
+ total_conditions = 0;
+ }
+ i--;
+ }
+ }
+ else {
+ i --;
+ }
+ }
+ }
+ }
+
+
+ if(computed) {
+ cudaMemcpy(results, results_cuda, sizeof(int) * steps / THREAD_NUM, cudaMemcpyDeviceToHost);
+
+ for(int j = 0; j < steps / THREAD_NUM; j++) {
+ total += results[j];
+ }
+
+ computed = false;
+ }
+
+ cudaMemcpy(masks_cuda, total_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+ cudaMemcpy(l_masks_cuda, total_l_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+ cudaMemcpy(r_masks_cuda, total_r_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+
+ solve_nqueen_cuda_kernel<<<steps/THREAD_NUM, THREAD_NUM>>>(n, n - mark, masks_cuda, l_masks_cuda, r_masks_cuda, results_cuda, total_conditions);
+
+ cudaMemcpy(results, results_cuda, sizeof(int) * steps / THREAD_NUM, cudaMemcpyDeviceToHost);
+
+ for(int j = 0; j < steps / THREAD_NUM; j++) {
+ total += results[j];
+ }
+
+ total *= 2;
+
+ if(n % 2 == 1) {
+ computed = false;
+ total_conditions = 0;
+
+ index = (1 << (n - 1) / 2);
+ m[0] |= index;
+
+ mask[1] = index;
+ l_mask[1] = index << 1;
+ r_mask[1] = index >> 1;
+ m[1] = (mask[1] | l_mask[1] | r_mask[1]);
+ i = 1;
+
+ while(i > 0) {
+ if((m[i] & t_mask) == t_mask) {
+ i--;
+ }
+ else {
+ index = (m[i] + 1) & ~m[i];
+ m[i] |= index;
+ if((index & t_mask) != 0) {
+ mask[i + 1] = mask[i] | index;
+ l_mask[i + 1] = (l_mask[i] | index) << 1;
+ r_mask[i + 1] = (r_mask[i] | index) >> 1;
+ m[i + 1] = (mask[i + 1] | l_mask[i + 1] | r_mask[i + 1]);
+ i++;
+ if(i == mark) {
+ total_masks[total_conditions] = mask[i];
+ total_l_masks[total_conditions] = l_mask[i];
+ total_r_masks[total_conditions] = r_mask[i];
+ total_conditions++;
+ if(total_conditions == steps) {
+ if(computed) {
+ cudaMemcpy(results, results_cuda, sizeof(int) * steps / THREAD_NUM, cudaMemcpyDeviceToHost);
+
+ for(int j = 0; j < steps / THREAD_NUM; j++) {
+ total += results[j];
+ }
+
+ computed = false;
+ }
+
+ // start computation
+ cudaMemcpy(masks_cuda, total_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+ cudaMemcpy(l_masks_cuda, total_l_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+ cudaMemcpy(r_masks_cuda, total_r_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+
+ solve_nqueen_cuda_kernel<<<steps/THREAD_NUM, THREAD_NUM>>>(n, n - mark, masks_cuda, l_masks_cuda, r_masks_cuda, results_cuda, total_conditions);
+
+ computed = true;
+
+ total_conditions = 0;
+ }
+ i--;
+ }
+ }
+ else {
+ i --;
+ }
+ }
+ }
+
+ if(computed) {
+ cudaMemcpy(results, results_cuda, sizeof(int) * steps / THREAD_NUM, cudaMemcpyDeviceToHost);
+
+ for(int j = 0; j < steps / THREAD_NUM; j++) {
+ total += results[j];
+ }
+
+ computed = false;
+ }
+
+ cudaMemcpy(masks_cuda, total_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+ cudaMemcpy(l_masks_cuda, total_l_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+ cudaMemcpy(r_masks_cuda, total_r_masks, sizeof(int) * total_conditions, cudaMemcpyHostToDevice);
+
+ solve_nqueen_cuda_kernel<<<steps/THREAD_NUM, THREAD_NUM>>>(n, n - mark, masks_cuda, l_masks_cuda, r_masks_cuda, results_cuda, total_conditions);
+
+ cudaMemcpy(results, results_cuda, sizeof(int) * steps / THREAD_NUM, cudaMemcpyDeviceToHost);
+
+ for(int j = 0; j < steps / THREAD_NUM; j++) {
+ total += results[j];
+ }
+ }
+
+ cudaFree(masks_cuda);
+ cudaFree(l_masks_cuda);
+ cudaFree(r_masks_cuda);
+ cudaFree(results_cuda);
+
+ delete[] total_masks;
+ delete[] total_l_masks;
+ delete[] total_r_masks;
+ delete[] results;
+
+ bunk = 1;
+
+ return total;
+}
+
+
+bool InitCUDA()
+{
+ int count;
+
+ cudaGetDeviceCount(&count);
+ if(count == 0) {
+ fprintf(stderr, "There is no device.\n");
+ return false;
+ }
+
+ int i;
+ for(i = 0; i < count; i++) {
+ cudaDeviceProp prop;
+ if(cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
+ if(prop.major >= 1) {
+ break;
+ }
+ }
+ }
+
+ if(i == count) {
+ fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
+ return false;
+ }
+
+ cudaSetDevice(i);
+
+ return true;
+}
+
+
+int main(int argc, char** argv)
+{
+ unsigned int hTimer;
+ double gpuTime;
+ // initialise card and timer
+ int deviceCount;
+ CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceCount(&deviceCount));
+ if (deviceCount == 0) {
+ fprintf(stderr, "There is no device.\n");
+ exit(EXIT_FAILURE);
+ }
+ int dev;
+ for (dev = 0; dev < deviceCount; ++dev) {
+ cudaDeviceProp deviceProp;
+ CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceProperties(&deviceProp, dev));
+ if (deviceProp.major >= 1)
+ break;
+ }
+ if (dev == deviceCount) {
+ fprintf(stderr, "There is no device supporting CUDA.\n");
+ exit(EXIT_FAILURE);
+ }
+ else
+ CUDA_SAFE_CALL(cudaSetDevice(dev));
+ CUT_SAFE_CALL( cutCreateTimer(&hTimer) );
+
+ int n = 8;
+ clock_t start, end;
+ long long solution;
+ bool cpu = true, gpu = true;
+ int argstart = 1, steps = 24576;
+
+ if(argc >= 2 && argv[1][0] == '-') {
+ if(argv[1][1] == 'c' || argv[1][1] == 'C') {
+ gpu = false;
+ }
+ else if(argv[1][1] == 'g' || argv[1][1] == 'G') {
+ cpu = false;
+ }
+
+ argstart = 2;
+ }
+
+ if(argc < argstart + 1) {
+ printf("Usage: %s [-c|-g] n steps\n", argv[0]);
+ printf(" -c: CPU only\n");
+ printf(" -g: GPU only\n");
+ printf(" n: n-queen\n");
+ printf(" steps: step for GPU\n");
+ printf("Default to 8 queen\n");
+ }
+ else {
+ n = atoi(argv[argstart]);
+ if(n <= 1 || n > 32) {
+ printf("Invalid n, n should be > 1 and <= 32\n");
+ printf("Note: n > 18 will require a very very long time to compute!\n");
+ return 0;
+ }
+
+ if(argc >= argstart + 2) {
+ steps = atoi(argv[argstart + 1]);
+ if(steps <= THREAD_NUM || steps % THREAD_NUM != 0) {
+ printf("Invalid step, step should be multiple of %d\n", THREAD_NUM);
+ return 0;
+ }
+ }
+ }
+
+ if(gpu) {
+ if(!InitCUDA()) {
+ return 0;
+ }
+
+ printf("CUDA initialized.\n");
+ }
+
+ if(cpu) {
+ CUDA_SAFE_CALL( cudaThreadSynchronize() );
+ CUT_SAFE_CALL( cutResetTimer(hTimer) );
+ CUT_SAFE_CALL( cutStartTimer(hTimer) );
+
+ //start = clock();
+ solution = solve_nqueen(n); //solve_nqueen_mcpu(n);
+ //solution = solve_nqueen(n);
+ //end = clock();
+ CUT_SAFE_CALL( cutStopTimer(hTimer) );
+ gpuTime = cutGetTimerValue(hTimer);
+
+ printf("CPU: %d queen = %lld time = %f msec\n", n, solution, gpuTime);
+ }
+
+ if(gpu) {
+ //start = clock();
+ CUDA_SAFE_CALL( cudaThreadSynchronize() );
+ CUT_SAFE_CALL( cutResetTimer(hTimer) );
+ CUT_SAFE_CALL( cutStartTimer(hTimer) );
+ solution = solve_nqueen_cuda(n, steps);
+ //end = clock();
+ CUT_SAFE_CALL( cutStopTimer(hTimer) );
+ gpuTime = cutGetTimerValue(hTimer);
+ printf("GPU: %d queen = %lld time = %f msec\n", n, solution, gpuTime);
+ }
+
+ return 0;
+}
|