From 11b308e7363e937966b035b4891db32b4eece3bf Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Fri, 1 Oct 2010 08:55:28 -0800 Subject: integrating recent changes from fermi-test into fermi (i'll use "fermi" for more disruptive changes to the pipeline model such as updating the MSHRs and getting rid of the warp tracker, ripping out DWF, etc...) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7805] --- benchmarks/CUDA/NQU/Makefile | 50 --- benchmarks/CUDA/NQU/README.GPGPU-Sim | 2 - benchmarks/CUDA/NQU/nqueen.cu | 758 ----------------------------------- 3 files changed, 810 deletions(-) delete mode 100644 benchmarks/CUDA/NQU/Makefile delete mode 100644 benchmarks/CUDA/NQU/README.GPGPU-Sim delete mode 100644 benchmarks/CUDA/NQU/nqueen.cu (limited to 'benchmarks/CUDA/NQU') diff --git a/benchmarks/CUDA/NQU/Makefile b/benchmarks/CUDA/NQU/Makefile deleted file mode 100644 index 35d46db..0000000 --- a/benchmarks/CUDA/NQU/Makefile +++ /dev/null @@ -1,50 +0,0 @@ -################################################################################ -# -# 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 deleted file mode 100644 index 978b6d5..0000000 --- a/benchmarks/CUDA/NQU/README.GPGPU-Sim +++ /dev/null @@ -1,2 +0,0 @@ -make -./gpgpu_ptx_sim__nqueen diff --git a/benchmarks/CUDA/NQU/nqueen.cu b/benchmarks/CUDA/NQU/nqueen.cu deleted file mode 100644 index 68dc9b5..0000000 --- a/benchmarks/CUDA/NQU/nqueen.cu +++ /dev/null @@ -1,758 +0,0 @@ -// 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