Changes

/**

* An Opimistic Optimistic approach to solving a Sudoku on a CUDA enabled GPU * Assumes that the puzzle is deterministic(single solvable solution) * AND each next step can be found with the kernel

* squares that have only a single appropiate value * OR values that (in the box) can only fit in one square

#include <stdio.h> #include <iostream> #include <cstdlib> #include <ctime> #include <iomanip> // CUDA header file #include "cuda_runtime.h" #include <device_launch_parameters.h> #ifndef __CUDACC__ #define __CUDACC__ #endif #include <device_functions.h> #include <stdio.h> // UNASSIGNED is used for empty cells in sudoku grid #define UNASSIGNED 0 // N is used for the size of Sudoku grid. Size will be NxN #define BOXWIDTH 5 #define N (BOXWIDTH * BOXWIDTH)

/* * kernel to solve a sudoku * Input: sudoku puzzle partitioned into boxes * * d_a = the sudoku puzzle * figures out what values can fit in each square * figures out how many spots each value can go * assigns the appropiate values, * saves to addedIdx to show that there is a change */__global__ void educatedGuess(int * d_a, int * addedIdx) { int idx = threadIdx.x + BOXWIDTH * threadIdx.y; int gridX = threadIdx.x + BOXWIDTH * blockIdx.x; int gridY = threadIdx.y + BOXWIDTH * blockIdx.y; int gridIdx = gridX + N * gridY; __shared__ bool hasValue[N]; //If the value occurs in the box __shared__ int inBox[N]; //Number of places each integer can go in the box hasValue[idx] = false; inBox[idx] = 0; __syncthreads(); int at = d_a[gridIdx]; if (at != 0) hasValue[at - 1] = true;

__global__ void educatedGuess(int * d_a, int * addedIdx) { int idx = threadIdx.x + BOXWIDTH * threadIdx.y; int gridX = threadIdx.x + BOXWIDTH * blockIdx.x; int gridY = threadIdx.y + BOXWIDTH * blockIdx.y; int gridIdx = gridX + N * gridY; __shared__ bool hasValue[N]; //If the value occurs in the box __shared__ int inBox[N]; //Number of places each integer can go in the box hasValue[idx] = false; inBox[idx] = 0; __syncthreads(); int at = d_a[gridIdx]; if (at != 0) hasValue[at - 1] = true; __syncthreads(); if (at != 0) return; //For remembering which values were seen in the rows and columns bool foundVal[N]; for (int i = 0; i < N; ++i) foundVal[i] = hasValue[i]; for (int check = 0; check < N; check++) { foundVal[d_a[N * check + gridX] - 1] = true; foundVal[d_a[N * gridY + check] - 1] = true; } int fndVals = 0; for( int i = 0; i < N; ++i) if (!foundVal[i]) { fndVals++; at = i + 1; } if (fndVals == 1) { //Only one possible value for this index d_a[gridIdx] = at; //assign value addedIdx[0] = gridIdx; //to tell host that the table has changed inBox[at - 1] = 4; //Prevent one index per value } __syncthreads(); //Calculate the number of places each integer can go in the box for (int i = 0; i < N; ++i) { int num = (idx + i) % N; //keep each thread on a seperate idx if (!foundVal[num]) inBox[num]++; __syncthreads(); } for (int i = 0; i < N; ++i) { //if there is only one possible index for that value assign the value if (inBox[i] == 1 && !foundVal[i]) { d_a[gridIdx] = i + 1; //assign value addedIdx[0] = gridIdx; //to tell host that the table has changed } } }

for /* Solves the Sudoku, with best values */ void SolveSudoku(int check = 0; check < grid[N][N; check++], int* d_a, int* d_results) { foundVal[d_a[N * check + gridX] dim3 block(BOXWIDTH, BOXWIDTH); int lastIdx(-1), nextIdx(- 1] ); do { lastIdx = truenextIdx; foundVal[ educatedGuess << <block, block >> > (d_a[N * gridY + check] - 1] , d_results); cudaMemcpy(&nextIdx, d_results, sizeof(int), cudaMemcpyDeviceToHost); } while (lastIdx != truenextIdx); }

int fndVals = 0; for( int i = 0; i < N; ++i) if (!foundVal[i]) { fndVals++; at = i + 1; } if (fndVals == 1) { /* A utility function to print grid */Only one possible value for this index d_a[gridIdx] = at; //assign value addedIdxvoid printGrid(int grid[0N] = gridIdx; //to tell host that the table has changed inBox[at - 1N] = 4; //Prevent one index per value } __syncthreads(); //Calculate the number of places each integer can go in the box { for (int i row = 0; i row < N; row++i) { int num = (idx + i) % N; //Ensure that each square is operating on a seperate idx if (!foundVal[num]) inBox[num]++; __syncthreads(); } for (int i col = 0; i col < N; col++i) { //if there is only one possible index for that value assign the value if printf(inBox"%3d", grid[irow] == 1 && !foundVal[icol]) { d_a[gridIdx] = i + 1; //assign value addedIdx[0] = gridIdx printf("\n"); //to tell host that the table has changed } } }

/* Takes a partially filled-in grid and attempts to assign values Driver Program totest above functions */ all unassigned locations in such a way to meet the requirements for Sudoku solution int main(non-duplication across rows, columns, and boxes) { /* 0 means unassigned cells */void SolveSudoku( int grid[N][N]= { {3, int* d_a0, 6, 5, 0, 8, 4, 0, 0}, int* d_results) {5, 2, 0, 0, 0, 0, 0, 0, 0}, dim3 block(BOXWIDTH {0, 8, 7, 0, 0, 0, 0, 3, 1}, BOXWIDTH); int lastIdx(- {0, 0, 3, 0, 1), nextIdx(-1);0, 0, 8, 0}, do {9, 0, 0, 8, 6, 3, 0, 0, 5}, lastIdx = nextIdx; {0, 5, 0, 0, 9, 0, 6, 0, 0}, educatedGuess << <block {1, 3, 0, 0, 0, 0, 2, block >> > (d_a5, 0}, d_results); cudaMemcpy(&nextIdx {0, 0, 0, 0, 0, 0, 0, d_results7, sizeof(int)4}, cudaMemcpyDeviceToHost); {0, 0, 5, 2, 0, 6, 3, 0, 0} } while (lastIdx != nextIdx);}

/* A utility function to print grid */void printGrid(int grid[N][N]){ for (int row = 0; row < N; row++) { for (int col = 0; col < N; col++) printf("%3d", grid[row][col]); printf("\n"); }} /* Driver Program to test above functions */int main(){ /* 0 means unassigned cells * int grid[N][N] = { {3, 0, 6, 5, 0, 8, 4, 0, 0}, {5, 2, 0, 0, 0, 0, 0, 0, 0}, {0, 8, 7, 0, 0, 0, 0, 3, 1}, {0, 0, 3, 0, 1, 0, 0, 8, 0}, {9, 0, 0, 8, 6, 3, 0, 0, 5}, {0, 5, 0, 0, 9, 0, 6, 0, 0}, {1, 3, 0, 0, 0, 0, 2, 5, 0}, {0, 0, 0, 0, 0, 0, 0, 7, 4}, {0, 0, 5, 2, 0, 6, 3, 0, 0} };  /** int grid[N][N] = {{0, 8, 0, 0, 0, 0, 0, 3, 0, 0, 0, 10, 9, 7, 11, 0}, {0, 9, 15, 13, 0, 10, 0, 0, 2, 6, 8, 16, 0, 0, 0, 0}, {0, 0, 16, 0, 15, 0, 8, 0, 9, 0, 0, 0, 6, 0, 2, 0}, {1, 0, 2, 0, 9, 11, 4, 6, 15, 3, 5, 7, 0, 0, 12, 0}, {16, 6, 4, 0, 5, 2, 0, 0, 1, 0, 0, 0, 11, 0, 0, 12}, {5, 11, 0, 0, 0, 3, 0, 15, 0, 16, 0, 13, 0, 1, 0, 8}, {0, 0, 3, 0, 0, 6, 11, 14, 0, 5, 7, 0, 0, 9, 0, 0}, {0, 0, 0, 14, 8, 0, 10, 0, 0, 11, 12, 0, 0, 0, 0, 0}, {0, 7, 13, 0, 0, 0, 0, 12, 0, 8, 9, 0, 0, 0, 3, 0}, {0, 0, 11, 9, 0, 7, 0, 0, 0, 0, 0, 12, 0, 8, 16, 5}, {0, 0, 10, 0, 11, 13, 0, 0, 0, 0, 0, 3, 12, 0, 6, 0}, {0, 5, 0, 0, 10, 15, 0, 1, 7, 2, 0, 0, 14, 11, 0, 0}, {0, 0, 5, 0, 0, 12, 14, 0, 0, 10, 0, 0, 15, 0, 0, 4}, {9, 0, 14, 6, 0, 0, 1, 0, 16, 0, 2, 0, 3, 0, 13, 0}, {8, 13, 0, 4, 0, 0, 0, 0, 12, 7, 3, 0, 0, 6, 0, 0}, {0, 16, 12, 0, 0, 5, 0, 9, 0, 13, 14, 4, 1, 0, 0, 0} }; /**/ int grid[N][N] = { {1, 0, 4, 0, 25, 0, 19, 0, 0, 10, 21, 8, 0, 14, 0, 6, 12, 9, 0, 0, 0, 0, 0, 0, 5},

int* d_a; //Table int* d_result; //Table change indicator  cudaMalloc((void**)&d_a, N*N * sizeof(int)); cudaMalloc((void**)&d_result, sizeof(int));

//Copy Sudoku over cudaMemcpy(d_a, grid, N*N * sizeof(int), cudaMemcpyHostToDevice);

SolveSudoku(grid, d_a, d_result);

//Copy Sudoku back cudaMemcpy(grid, d_a, N*N * sizeof(int), cudaMemcpyDeviceToHost); printGrid(grid);

cudaFree(d_a); cudaFree(d_result);

return 0;