From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Thu, 15 Jul 2010 18:09:46 -0800 Subject: creating branch for adding support for CUDA 3.x and Fermi [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829] --- benchmarks/CUDA/NQU/Makefile | 50 +++ benchmarks/CUDA/NQU/README.GPGPU-Sim | 2 + benchmarks/CUDA/NQU/nqueen.cu | 758 +++++++++++++++++++++++++++++++++++ 3 files changed, 810 insertions(+) create mode 100644 benchmarks/CUDA/NQU/Makefile create mode 100644 benchmarks/CUDA/NQU/README.GPGPU-Sim create mode 100644 benchmarks/CUDA/NQU/nqueen.cu (limited to 'benchmarks/CUDA/NQU') diff --git a/benchmarks/CUDA/NQU/Makefile b/benchmarks/CUDA/NQU/Makefile new file mode 100644 index 0000000..35d46db --- /dev/null +++ b/benchmarks/CUDA/NQU/Makefile @@ -0,0 +1,50 @@ +################################################################################ +# +# Copyright 1993-2006 NVIDIA Corporation. All rights reserved. +# +# NOTICE TO USER: +# +# This source code is subject to NVIDIA ownership rights under U.S. and +# international Copyright laws. +# +# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE +# CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR +# IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH +# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF +# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. +# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, +# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS +# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE +# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE +# OR PERFORMANCE OF THIS SOURCE CODE. +# +# U.S. Government End Users. This source code is a "commercial item" as +# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of +# "commercial computer software" and "commercial computer software +# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) +# and is provided to the U.S. Government only as a commercial end item. +# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through +# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the +# source code with only those rights set forth herein. +# +################################################################################ +# +# Build script for project +# +################################################################################ + +# Add source files here +EXECUTABLE := nqueen +# Cuda source files (compiled with cudacc) +CUFILES := nqueen.cu +# C/C++ source files (compiled with gcc / c++) +CCFILES := + +GPGPUSIM_ROOT := ../../.. + +################################################################################ +# Rules and targets + +include ../../../common/common.mk + + diff --git a/benchmarks/CUDA/NQU/README.GPGPU-Sim b/benchmarks/CUDA/NQU/README.GPGPU-Sim new file mode 100644 index 0000000..978b6d5 --- /dev/null +++ b/benchmarks/CUDA/NQU/README.GPGPU-Sim @@ -0,0 +1,2 @@ +make +./gpgpu_ptx_sim__nqueen 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 +#include +#include + +#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<<>>(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<<>>(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<<>>(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<<>>(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; +} -- cgit v1.3