Changes

// CUDA header file #include "cuda_runtime.h" #include <device_launch_parameters.h> #ifndef __CUDACC__ #define __CUDACC__ #endif // UNASSIGNED is used for empty cells in Sudoku grid #define UNASSIGNED 0 // BOX_W is used for the length of one of the square sub-regions of the Sudoku grid. // Overall length will be N * N. #define BOX_W 5 #define N (BOX_W * BOX_W) __global__ void solve(int* d_a) { // Used to remember which row | col | box ( section ) have which values __shared__ bool rowHas[N][N]; __shared__ bool colHas[N][N]; __shared__ bool boxHas[N][N];

// Used to ensure that BOX_W is used for the table has changedlength of one of the square sub-regions of the Sudoku grid.// Overall length will be N * N.#define BOX_W 5 __shared__ bool changed;#define N (BOX_W * BOX_W)

__global__ void solve(int* d_a) { // Where the square is located in the SudokuUsed to remember which row | col | box ( section ) have which values int row = threadIdx.x __shared__ bool rowHas[N][N]; int col = threadIdx.y __shared__ bool colHas[N][N]; int box = row / BOX_W + (col / BOX_W) * BOX_W __shared__ bool boxHas[N][N];

// Unique identifier for each square in row, col, box // Corresponds Used to ensure that the generic Sudoku Solve // Using a Sudoku to solve a Sudoku !!!table has changed int offset = col + (row % BOX_W) * BOX_W + (box % BOX_W) __shared__ bool changed;

// Square's location Number of spaces which can place the number in the Sudokueach section __shared__ int rowCount[N][N]; __shared__ int colCount[N][N]; __shared__ int gridIdx = col * boxCount[N][N + row];

// Where the square is located in the Sudoku int at row = d_a[gridIdx]threadIdx.x; int col = threadIdx.y; int box = row / BOX_W + (col / BOX_W) * BOX_W;

bool notSeen[N]; // Unique identifier for each square in row, col, box for ( // Corresponds to the generic Sudoku Solve // Using a Sudoku to solve a Sudoku !!! int i offset = 0; i < N; col +(row % BOX_W) * BOX_W +i(box % BOX_W) notSeen[i] = true;

rowHas[col][row] = false; // Square's location in the Sudoku colHas[col][row] int gridIdx = false; boxHas[col][* N + row] = false; __syncthreads();

if ( int at != UNASSIGNED) { rowHasd_a[rowgridIdx][at - 1] = true; colHas[col][at - 1] = true; boxHas[box][at - 1] = true; } // Previous loop has not changed any values do { // RESET counters rowCount[col][row] = 0; colCount[col][row] = 0; boxCount[col][row] = 0; __syncthreads();

if (gridIdx == 0) // forget previous change changed = false bool notSeen[N]; int count = 0; // number of values which can fit in this square int guess = 0; // last value found which can fit in this square for (int idx i = 0; idx i < N; ++idx) { // Ensures that every square in each section is working on a different number in the section int num = (idx + offset) % N; if (at == UNASSIGNED && notSeen[num]) { if (rowHas[row][num] || boxHas[box][num] || colHas[col][num]i) notSeen[numi] = false; else { ++count; guess = num; rowCount[row][num]++; colCount[col][num]++true; boxCount[box][num]++; } } __syncthreads(); }

// Find values which can go in only one spot in the section for (int idx = 0; idx < N && count > 1; ++idx) { if (notSeen rowHas[idxcol] && (rowCount[row][idx] == 1 || boxCountfalse; colHas[boxcol][idxrow] == 1 || colCountfalse; boxHas[col][idxrow] == 1)) { // In this section this value can only appear in this square guess = idxfalse; count = 1 __syncthreads(); } }

if (count =at != 1UNASSIGNED) { at = guess + 1; rowHas[row][guessat - 1] = true; colHas[col][guessat - 1] = true; boxHas[box][guessat - 1] = true; } // Previous loop has not changed = true;any values } do { __syncthreads(); // RESET counters } while (changed) rowCount[col][row] = 0; d_a colCount[col][gridIdxrow] = at0; } bool validate(int result[N boxCount[col][Nrow]= 0; __syncthreads() {;  for if (int row gridIdx == 0) // forget previous change changed = false; int count = 0; row < N; row++) // number of values which can fit in this square int guess = 0; // last value found which can fit in this square for (int col idx = 0; col idx < N; col++idx) { // Ensures that every square in each section is working on a different number in the section int num = (idx + offset)% N; if (resultat == UNASSIGNED && notSeen[num]) { if (rowHas[row][colnum] || boxHas[box][num] || colHas[col][num] == 0) return notSeen[num] = false; return true else { ++count; } guess = num; void print(int result[N rowCount[row][Nnum]) { for (int row = 0; row < N; row++) {; for (int col = 0; col < N; col colCount[col][num]++); printf("%3d", result boxCount[rowbox][colnum])++; printf } } __syncthreads("\n"); } } // Driver program to test main program functionsFind values which can go in only one spot in the section int main for () { int h_aidx = 0; idx < N && count > 1; ++idx) { if (notSeen[Nidx] && (rowCount[row][Nidx] = { { = 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},|| boxCount[box][idx] == 1 || colCount[col][idx] == 1)) { // In this section this value can only appear in this square guess = idx; count = 1; } }  if (count == 1) { at = guess + 1; { 5, 0, 19, 23, 24, 0, 22, 12, 0, 0, 16, 6, 0, 20, 0, 18, 0, 25, 14, 13, 10, 11, 0, 1, 15 rowHas[row][guess] = true; colHas[col][guess] = true; boxHas[box][guess] = true; changed = true; }, { 0, 0, 0, 0, 0, 0, 21, 5, 0, 20, 11, 10, 0, 1, 0, 4, 8, 24, 23, 15, 18, __syncthreads(); } while (changed);  //SOLVED CHECK if (!(rowHas[row][col] || colHas[row][col] || boxHas[row][col])) changed = true; __syncthreads(); if (changed && gridIdx == 0, 16, 22, 19) at = 0;  d_a[gridIdx] = at;},  void print(int result[N][N]) { for (int row = 0, 7, 21, 8, 18, ; row < N; row++) { for (int col = 0; col < N; col++) printf("%3d", 0, 0, 11, 0, 5, 0, 0, 24, 0, 0, 0, 17, 22, result[row][col]); printf("\n"); }} // Driver program to test main program functionsint main() { int h_a[N][N] = { { 1, 90, 4, 60, 25, 0, 19, 0}, { 0, 1310, 1521, 8, 0, 22, 14, 0, 18, 06, 1612, 09, 0, 0, 40, 0, 0, 0, 19 5}, { 05, 0, 019, 23, 24, 20 0, 21, 17}22, { 12, 0, 11, 0, 16, 6, 0, 20, 0, 18, 0, 025, 14, 1513, 010, 011, 0, 01, 2115}, { 0, 25 0, 19 0, 0, 40, 0, 2221, 14 5, 0, 20, 11, 10, 0}, { 81, 0, 04, 8, 24, 23, 15, 2118, 0, 16, 022, 19}, { 0, 07, 221, 08, 318, 0, 0, 0, 11, 0, 17, 23, 18, 22, 05, 0, 0, 24, 6}0, { 40, 0, 1417, 1822, 71, 9, 06, 2225, 21 0, 19 0}, { 0, 13, 15, 0, 22, 14, 0, 218, 0, 516, 0, 0, 0, 64, 16 0, 15 0, 0, 1119, 12} 0, { 22 0, 0, 24, 020, 21, 2317}, 0 { 12, 0, 11, 0, 76, 0, 0, 40, 0, 1415, 0, 20, 12 0, 0, 821, 525, 19, 0, 25 4, 9}0, 22, { 2014, 0, 020, 0}, { 58, 0, 0, 21, 0, 16, 0, 0, 17 0, 92, 0, 12, 18 3, 0, 10, 0, 0, 717, 23, 18, 2422, 0, 0, 0, 1324, 46}, { 13 4, 0, 014, 518, 07, 2, 23, 149, 4, 180, 22, 021, 1719, 0, 0, 20, 0, 12, 90, 21 5, 12 0, 0, 0, 86, 16, 15, 0, 11, 12}, { 14, 2322, 0, 24, 0, 23, 0, 0, 11, 0, 07, 0, 0, 4, 0, 14, 0, 20 2, 2512, 0, 38, 45, 1319, 0, 11, 2125, 9, 5, 18, 22}, { 20, 70, 0, 0, 11 5, 17, 20, 24 0, 0, 0, 0, 317, 49, 10, 12, 18, 0, 01, 6, 140, 0, 57, 2524, 13 0, 0, 0, 13, 04}, { 13, 0, 0, 16 5, 90, 02, 1723, 1114, 74, 1018, 2522, 0, 17, 0, 0, 1320, 60, 01, 09, 1821, 012, 0, 19, 4, 0, 08, 2011}, { 614, 1523, 0, 1924, 4, 130, 0, 0, 50, 0, 18, 11, 0, 0, 90, 8, 22, 1620, 25, 10, 70, 03, 04, 013, 0}, { 11, 21, 09, 05, 18, 022}, { 27, 0, 0, 1011, 17, 1920, 324, 0, 10, 0, 22, 93, 4, 11, 15, 01, 2012, 0, 0, 86, 2314, 0, 25}, { 05, 24, 825, 13, 10, 0, 0}, { 4, 200, 0, 1716, 14 9, 0, 017, 1811, 07, 1610, 2225, 50, 0, 11, 0, 1013, 06, 0}, { 23 0, 1018, 0, 0, 19, 04, 0, 0, 20}, { 06, 1815, 0, 619, 04, 1613, 0, 0, 17, 15, 0, 1318, 11, 0, 0, 39, 19 8, 1222, 0}16, { 25, 10, 57, 0, 14, 11, 0, 17, 0, 80}, 24 { 0, 13 0, 0, 19, 23, 15, 92, 0, 0, 1210, 19, 3, 0, 20 1, 0, 22, 09, 7}4, 11, { 015, 0, 1720, 40, 0, 22 8, 1523, 0, 23, 11, 12, 25}, { 0, 024, 08, 13, 0, 181, 80, 0, 74, 20, 0, 017, 14, 0, 13} 0, { 1918, 60, 2316, 22, 85, 0, 11, 0, 1, 2510, 4, 14, 20, 0}, 3, 7, 13 { 23, 10, 11, 16, 0, 0, 0, 0, 0, 0}, { 18, 0, 46, 0, 1716, 0, 3, 0, 2417, 01, 80, 20, 23, 11, 10, 25, 2213, 0, 0, 03, 19, 12, 13, 2, 18, 6, 0}, { 25, 05, 0, 714, 1611, 0, 17, 0, 68, 24, 1713, 20, 2119, 023, 1815, 09, 0, 0, 1912, 0, 20, 0, 822, 0, 07}, { 0, 0, 17, 4, 0}, { 1822, 15, 90, 2523, 111, 212, 1125, 0, 0, 13, 22, 40, 0, 2118, 0, 58, 0, 23, 7, 0, 0, 1514, 0, 3, 0, 813}, { 19, 06, 2123, 1022, 08, 0, 12, 0, 20 1, 1625, 04, 1914, 2, 0, 3, 7, 13, 10, 11, 16, 0, 0, 0, 0, 0, 15, 14, 4, 2, 18, 23, 25, 11, 7, 0} 0, 0, 0}, { 0, 4, 0, 17, 0, 3, 0, 24, 0, 8, 20, 23, 11, 10, 25, 22, 0, 0, 0, 12, 13, 2, 18, 6, 0}, { 0, 0, 7, 16, 0, 0, 6, 17, 2, 21, 0, 18, 0, 0, 0, 19, 0, 0, 8, 0, 0, 0, 0, 4, 0}, { 18, 9, 25, 1, 2, 11, 0, 0, 13, 22, 4, 0, 21, 0, 5, 0, 23, 7, 0, 0, 15, 0, 3, 0, 8}, { 0, 21, 10, 0, 0, 12, 0, 20, 16, 0, 19, 0, 0, 0, 0, 15, 14, 4, 2, 18, 23, 25, 11, 7, 0} };  int* d_a; //Table cudaMalloc((void**)&d_a, N * N * sizeof(int)); // Copy Sudoku to device cudaMemcpy(d_a, h_a, N * N * sizeof(int), cudaMemcpyHostToDevice); dim3 dBlock(N, N); solve << <1, dBlock >> > (d_a); // Copy Sudoku back to host cudaMemcpy(h_a, d_a, N * N * sizeof(int), cudaMemcpyDeviceToHost); // Check if solved if (h_a[0][0]) print(h_a); else printf("No solution could be found."); cudaFree(d_a); return 0; };