112
edits
Changes
→Kernel Optimization Attempts
Sudoku Solver Profiling
Rather than try to continuously increase the difficulty of a 9x9 sudoku, I decided to modify the program I found to handle larger and large sudokus, increasing the size of the matrices that make up the sudoku (starting with a 9x9 sudoku, which is 9 3x3 matrices, then 16x16 which is 16 4x4 matrices, and finally 25x25 which is 25 5x5 matrices) without changing the logic of the program (only constants), so larger sudokus are solved the same way as a normal one.
Source code from: https://www.geeksforgeeks.org/sudoku-backtracking-7/
{| class="wikitable mw-collapsible mw-collapsed"
! Original Code:
|-
|
// A Backtracking program in C++ to solve Sudoku problem
/* Check if 'num' is not already placed in current row,
current column and current 3x3 box */
return !UsedInRow(grid, row, num) && !UsedInCol(grid, col, num) && !UsedInBox(grid, row - row%3 , col - col%3, num)&& grid[row][col]==UNASSIGNED;
}
/* Check if 'num' is not already placed in current row,
current column and current 4x4 box */
return !UsedInRow(grid, row, num) && !UsedInCol(grid, col, num) && !UsedInBox(grid, row - row%4 , col - col%4, num)&& grid[row][col]==UNASSIGNED;
}
/* Check if 'num' is not already placed in current row,
current column and current 5x5 box */
return !UsedInRow(grid, row, num) && !UsedInCol(grid, col, num) && !UsedInBox(grid, row - row%5 , col - col%5, num)&& grid[row][col]==UNASSIGNED;
}
return 0;
}
|}
Obtaining flat profiles and call graphs on matrix environment:
Attempted to run the program with a number of files (8K resolution):
{| class="wikitable mw-collapsible mw-collapsed"
! Sample Images
|-
|
[[File:Cabin small.jpg]]
[[File:Cabin2 small.jpg]]
|}
{| class="wikitable mw-collapsible mw-collapsed"
[[File:Julia.jpg]]
|}
This problem would be fairly simple to parallelize. In the image created by Julia sets each pixel is independent of the others. This problem involves Complex numbers, but that can be simply represented by using two arrays, or pairs of floats.
==== Assignment 1: Selection for parallelizing ====
After reviewing the three programs above, we decided to attempt to parallelize the Sudoku Solver Program for a few reasons.
1. By increasing the dimensions of the smaller matrices that make up a sudoku by one, we see a major increase in the time it takes to solve the sudoku, from almost instantly to around 38 seconds, and then to '''36 minutes'''. With a 25x25 sudoku (of 5x5 matrices), several functions were called over '''100 million times'''.
2. Based on the massive time increases and similarity to the Hamiltonian Path Problem [https://www.hackerearth.com/practice/algorithms/graphs/hamiltonian-path/tutorial/] which also uses backtracking to find a solution, we believe the run time of the sudoku solver to have a Big O notation that approaches O(n!) where 'n' is the number of blank spaces in the sudoku as the sudoku solver uses recursion to check every single possible solution, returning to previous steps if the tried solution does not work. O(n!) is an even worse runtime than O(n^2).
3. The Julia sets still took less than 6 minutes after increasing the image size, and the EasyBMP only took a few seconds to convert a large, high resolution image. Therefore, the Sudoku Solver had the greatest amount of time to be shaven off through optimization and thus offered the most challenge.
=== Assignment 2 ===
This code is unable to solve the 16x16 in any reasonable amount of time (I stopped it at 10+ minutes).
If you consider the 130+ empty spaces in the grid I estimate over 130^2 calls to cudaMemcpy either way...
So we need an algorithm which will check each open spot, calculate all possible values which can fit there, and assign single values.
We can also check each section (Box, row, col) for values which can only go in one place
{| class="wikitable mw-collapsible mw-collapsed"
! Faster than YoursAttempt One...
|-
|
'''Single Pass Sudoku Solver'''
This Kernel was designed to run on a single block with dimensions N*N the size of the Sudoku
limiting us to a Sudoku of size 25 * 25
For each empty space, counts the number possible values which can fit and how many times each value can fit in that section
If only one value can fit or that value has only one place, assigns the value
__global__ void superSolve(int * d_a) {
}
}
[[File:Backtrack_vs_Kernel.png]]
=== Assignment 3 ===
Changes:
Reduced Thread Divergence/CGMA
-each thread now remembers which values it has seen in a boolean array
- values are only assigned to the grid after the kernel 'solves' the sudoku
- at value in kernel and shared memory for rowHas, colHas, boxHas, updated in a single place
Coalesced Memory
- change modifying _Has and _Count arrays from row->col to col->row as row(threadIdx.x) is our fastest moving dimension
Clarified Code
- use gridIdx == 0 rather then !gridIdx
- use a do-while loop rather then a while loop
{| class="wikitable mw-collapsible mw-collapsed"
! Optimized kernel : Single PassFull code
|-
|
|}
[[File:Unoptimized_vs_Optimized.png]] ===Kernel Optimization Attempts===These Kernels change a minor part of the Optimized Kernel or use a slightly different algorithm in an attempt to make it faster ChangesChange :Replaces the boolean array hasSeen with a single int & uses bitwise operators Used faster Theory : Since local array variables of threads are stored in Global memory this was an attempt to move that into a register Result : No speed up noticed, suggesting that more is happening beyond arrays stored in Global memory, perhaps some type of paging, more testing would be needed on something less erratic then a Sudoku Solver {| class="wikitable mw-collapsible mw-collapsed"! Using a int as a boolean array|-each thread now remembers | __global__ void solve(int* d_a) { // Used to remember which row | col | box ( section ) have which values it __shared__ bool rowHas[N][N]; __shared__ bool colHas[N][N]; __shared__ bool boxHas[N][N]; // Used to ensure that the table has seen changed __shared__ bool changed; // Number of spaces which can place the number in each section __shared__ int rowCount[N][N]; __shared__ int colCount[N][N]; __shared__ int boxCount[N][N]; // Where the square is located in the Sudoku int row = threadIdx.x; int col = threadIdx.y; int box = row / BOX_W + (col / BOX_W) * BOX_W; int gridIdx = col * N + row; int at = d_a[gridIdx]; // Unique identifier for each square in row, col, box // Corresponds to the generic Sudoku Solve // Using a Sudoku to solve a boolean arraySudoku !!! int offset = col + (row % BOX_W) * BOX_W + (box % BOX_W); // Square's location in the Sudoku Reduced Thread Divergence int notSeen = 0; rowHas[col][row] = false; colHas[col][row] = false; boxHas[col][row] = false; __syncthreads(); if (at != UNASSIGNED) { rowHas[row][at - 1] = true; colHas[col][at - 1] = true; boxHas[box][at - 1] = true; } else { notSeen = ~0; } __syncthreads(); // 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; int count = 0; // number of values which can fit in this square int guess = 0; // last value found which can fit in this square int b_shuttle = 1; for (int idx = 0; idx < N; ++idx) { // Ensures that every square in each section is working on a different number in the section int num = (idx + offset) % N; if (b_shuttle & notSeen) { if (rowHas[row][num] || boxHas[box][num] || colHas[col][num]) notSeen ^= b_shuttle; else { ++count; guess = num; rowCount[row][num]++; colCount[col][num]++; boxCount[box][num]++; } } b_shuttle <<= 1; __syncthreads(); } // Find values are which can go in only assigned to one spot in the grid after section b_shuttle = 1; for (int idx = 0; idx < N && count > 1; ++idx) { int num = (idx + offset) % N; if ((b_shuttle & notSeen) && (rowCount[row][num] == 1 || boxCount[box][num] == 1 || colCount[col][num] == 1)) { // In this section this value can only appear in this square guess = num; count = 1; } b_shuttle <<= 1; } if (count == 1) { at = guess + 1; notSeen = 0; rowHas[row][guess] = true; colHas[col][guess] = true; boxHas[box][guess] = true; changed = true; } __syncthreads(); } while (changed); //SOLVED CHECK if (!(rowHas[row][col] || colHas[row][col] || boxHas[row][col])) changed = true; __syncthreads(); if (changed && gridIdx == 0) at = 0; d_a[gridIdx] = at; }|} Change : Remove the kernel counters, and logic which checks for a section needing a value in one place Theory : The counting logic requires a additional nested loop each solve cycle and created more thread divergence Result : The algorithm is slower, probably because 'solvessections requiring a single value' adds more values early in the sudoku removing wait times kernel resulting in less passes overall Also this kernel is similar to one of my earlier builds, which was unable to solve the 9x9 getting stuck on every square having more then one possible value{| class="wikitable mw-collapsible mw-collapsed"! Dropping Section Logic|-| __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 the table has changed __shared__ bool changed; // Number of spaces which can place the number in each section // Where the square is located in the Sudoku int row = threadIdx.x; int col = threadIdx.y; int box = row / BOX_W + (col / BOX_W) * BOX_W; // Unique identifier for assigning each square in row, col, box // Corresponds to global memorythe generic Sudoku Solve // Using a Sudoku to solve a Sudoku !!! int offset = col + (row % BOX_W) * BOX_W + (box % BOX_W); // Square's location in the Sudoku int gridIdx = col * N + row; int at = d_a[gridIdx]; bool notSeen[N]; for (int i = 0; i < N; ++i) notSeen[i] = true; rowHas[col][row] = false; colHas[col][row] = false; boxHas[col][row] = false; __syncthreads(); if (at != UNASSIGNED) { rowHas[row][at - 1] = true; colHas[col][at - 1] = true; boxHas[box][at - 1] = true; } // Previous loop has not changed any values do { // RESET counters __syncthreads(); if (gridIdx == 0) // forget previous change changed = false; int count = 0; // number of values which can fit in this square int guess = 0; // last value found which can fit in kernel this square for (int idx = 0; idx < 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]) notSeen[num] = false; else { ++count; guess = num; } } __syncthreads(); } if (count == 1) { at = guess + 1; rowHas[row][guess] = true; colHas[col][guess] = true; boxHas[box][guess] = true; changed = true; } __syncthreads(); } while (changed); //SOLVED CHECK if (!(rowHas[row][col] || colHas[row][col] || boxHas[row][col])) changed = true; __syncthreads(); if (changed && gridIdx == 0) at = 0; d_a[gridIdx] = at; }|} Change : Quickly finds one section that requires a single value in one spot, by checking all sections at once and rowHasremembering a single section Theory : Similar to the previous Kernel, colHastrying to remove the second loop Result : Surprisingly slow, boxHasgains little benefit from the section logic and shared memory, updated in yet is still required to count all values{| class="wikitable mw-collapsible mw-collapsed"! Notify - Determines a single section that has a limited value (removes section loop)|-| __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 the table has changed __shared__ bool changed; // Number of spaces which can placethe number in each section __shared__ int rowCount[N][N]; __shared__ int colCount[N][N]; __shared__ int boxCount[N][N]; // Where the square is located in the Sudoku int row = threadIdx.x; int col = threadIdx.y; int box = row / BOX_W + (col / BOX_W) * BOX_W; // Unique identifier for each square in row, reducing wait time col, box // Corresponds to the generic Sudoku Solve // Using a Sudoku to solve a Sudoku !!! int offset = col + (row % BOX_W) * BOX_W + (box % BOX_W); // Square's location in the Sudoku int gridIdx = col * N + row; int at = d_a[gridIdx]; bool notSeen[N]; for updating(int i = 0; i < N; ++i) notSeen[i] = true; rowHas[col][row] = false; colHas[col][row] = false; boxHas[col][row] = false; Coalesced Memory __syncthreads(); __shared__ int notify; if (at != UNASSIGNED) { rowHas[row][at - change modifying _Has and _Count arrays from 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 to ] = 0; boxCount[col][row] = 0; __syncthreads(); if (gridIdx == 0) { // forget previous change changed = false; notify = -1; } 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 = 0; idx < 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]) notSeen[num] = false; else { ++count; guess = num; rowCount[row][num]++; colCount[col][num]++; boxCount[box][num]++; } } __syncthreads(); } if (rowCount[row][col] == 1 || colCount[row][col] == 1 || boxCount[row][col] == 1) notify = col; __syncthreads(); // Find values which can go in only one spot in the section if (notify >0 && at == UNASSIGNED && notSeen[notify] && (rowCount[row][notify] == 1 || boxCount[box][notify] == 1 || colCount[col][notify] == 1)) { // In this section this value can only appear in this square guess = notify; count = 1; } if (count == 1) { at = guess + 1; rowHas[row][guess] = true; colHas[col][guess] = true; boxHas[box][guess] = true; changed = true; } __syncthreads(); } while (changed); //SOLVED CHECK if (!(rowHas[row][col ] || colHas[row][col] || boxHas[row][col])) changed = true; __syncthreads(); if (changed && gridIdx == 0) at = 0; d_a[gridIdx] = at; }|} Change : Refactors the algorithm to count the total numbers that can fit in a square or section Then counts down as values are added Theory : Remove redundant counting logic that occurred during the Optimized Kernel each pass Result : Not faster, HOWEVER there is a slight error, by setting notSeen = 0, the section counters will rarely reach one{| class="wikitable mw-collapsible mw-collapsed"! CountDown - using Int as Boolean Array(EDITED now 4.28 seconds)|-| __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 the table has changed __shared__ bool changed; // Number of spaces which can place the number in each section __shared__ int rowCount[N][N]; __shared__ int colCount[N][N]; __shared__ int boxCount[N][N]; // Where the square is located in the Sudoku int row = threadIdx.x; int col = threadIdx.y; int box = row / BOX_W + (col / BOX_W) * BOX_W; int gridIdx = col * N + row; int at = d_a[gridIdx]; // Unique identifier for each square in row, col, box // Corresponds to the generic Sudoku Solve // Using a Sudoku to solve a Sudoku !!! int offset = col + (row % BOX_W) is our fastest moving dimension* BOX_W + (box % BOX_W); - query // Square's location in order the Sudoku int count = 0; //Number of values which can fit in this square int notSeen = 0; //Boolean Array as an Integer if (gridIdx == 0) changed = true; rowHas[col][row] = false; colHas[col][row] = false; boxHas[col][row] = false; rowCount[col][row] = 0; colCount[col][row] = 0; boxCount[col][row] = 0; __syncthreads(); if (at != UNASSIGNED) { rowHas[row][at ->1] = true; colHas[col][at - 1] = true; boxHas[box][at ->1] = true; } __syncthreads(); int guess; int b_shuttle = 1; for (int idx = 0; idx < N; ++idx) { int num = (idx + offset) % N; if (at == UNASSIGNED && !(rowHas[row][num] || boxHas[box][num] || colHas[col][num])) { notSeen |= b_shuttle; //this value can go here ++count; //how many values this square can have guess = num; //how many values this section can have rowCount[row][num]++; colCount[col][num]++; boxCount[box][num]++; } __syncthreads(); b_shuttle <<= 1; } if (at == UNASSIGNED && count == 0) //NOT POSSIBLE SUDOKU changed = false; __syncthreads(); if (count == 1) { at = guess + 1; notSeen = count = 0; rowHas[row][guess] = true; colHas[col ][guess] = true; boxHas[box][guess] = true; } // Previous loop has not changed any values while (changed) { __syncthreads(); if (gridIdx == 0) // forget previous change changed = false; int b_shuttle = 1; for (int idx = 0; idx < N; ++idx) { // Ensures that every square in each section is working on a different number in the same reasonsection int num = (idx + offset) % N; if (b_shuttle & notSeen && (at != UNASSIGNED || rowHas[row][num] || boxHas[box][num] || colHas[col][num])) { rowCount[row][num]--; colCount[col][num]--; boxCount[box][num]--; notSeen ^= b_shuttle; --count; } __syncthreads(); if (b_shuttle & notSeen && (count == 1 || rowCount[row][num] == 1 || boxCount[box][num] == 1 || colCount[col][num] == 1)) { rowHas[row][num] = true; colHas[col][num] = true; boxHas[box][num] = true; changed = true; notSeen ^= b_shuttle; at = num + 1; count = 0; } b_shuttle <<= 1; } __syncthreads(); }; if (!(rowHas[row][col] && colHas[row][col] && boxHas[box][col])) changed = true; //HAVE NOT SOLVED the sudoku __syncthreads(); if (changed && gridIdx == 0) at = 0; d_a[gridIdx] = at; Clarified Code} |} Change : uses countdown logic with a boolean array Result : Similar times to other Countdown kernel {| class="wikitable mw-collapsible mw-collapsed"! Countdown Boolean Array (EDITED - now 4.37ms)|- use | __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 the table has changed __shared__ bool changed; // Number of spaces which can place the number in each section __shared__ int rowCount[N][N]; __shared__ int colCount[N][N]; __shared__ int boxCount[N][N]; // Where the square is located in the Sudoku int row = threadIdx.x; int col = threadIdx.y; int box = row / BOX_W + (col / BOX_W) * BOX_W; int gridIdx = col * N + row; int at = d_a[gridIdx]; // Unique identifier for each square in row, col, box // Corresponds to the generic Sudoku Solve // Using a Sudoku to solve a Sudoku !!! int offset = col + (row % BOX_W) * BOX_W + (box % BOX_W); // Square's location in the Sudoku int count = 0; //Number of values which can fit in this square bool notSeen[N]; //Boolean Array as an Integer for(int idx = 0; idx < N; ++idx) notSeen[idx] = false; if (gridIdx == 0 rather then ) changed = true; rowHas[col][row] = false; colHas[col][row] = false; boxHas[col][row] = false; rowCount[col][row] = 0; colCount[col][row] = 0; boxCount[col][row] = 0; __syncthreads(); if (at !gridIdx= UNASSIGNED) { rowHas[row][at - 1] = true; colHas[col][at - use a do1] = true; boxHas[box][at -1] = true; } __syncthreads(); int guess; for (int idx = 0; idx < N; ++idx) { int num = (idx + offset) % N; if (at == UNASSIGNED && !(rowHas[row][num] || boxHas[box][num] || colHas[col][num])) { notSeen[num] = true; //this value can go here ++count; //how many values this square can have guess = num; //how many values this section can have rowCount[row][num]++; colCount[col][num]++; boxCount[box][num]++; } __syncthreads(); } if (at == UNASSIGNED && count == 0) //NOT POSSIBLE SUDOKU changed = false; __syncthreads(); if (count == 1) { at = guess + 1; count = 0; notSeen[guess] = false; rowHas[row][guess] = true; colHas[col][guess] = true; boxHas[box][guess] = true; } // Previous loop has not changed any values while loop rather then (changed) { __syncthreads(); if (gridIdx == 0) // forget previous change changed = false; for (int idx = 0; idx < N; ++idx) { // Ensures that every square in each section is working on a while loopdifferent number in the section int num = (idx + offset) % N; if (notSeen[num] && (at != UNASSIGNED || rowHas[row][num] || boxHas[box][num] || colHas[col][num])) { rowCount[row][num]--; colCount[col][num]--; boxCount[box][num]--; notSeen[num] = false; --count; } __syncthreads(); if ( notSeen[num] && (count == 1 || rowCount[row][num] == 1 || boxCount[box][num] == 1 || colCount[col][num] == 1)) { rowHas[row][num] = true; colHas[col][num] = true; boxHas[box][num] = true; changed = true; notSeen[num] = false; at = num + 1; count = 0; } } __syncthreads(); }; if (!(rowHas[row][col] && colHas[row][col] && boxHas[box][col])) changed = true; //HAVE NOT SOLVED the sudoku __syncthreads(); if (changed && gridIdx == 0) at = 0; d_a[gridIdx] = at; }|}[[File:Kernel_Compare.png]]==== Occupancy Calculations ===={| class="wikitable mw-collapsible mw-collapsed"! For 9x9:|-|[[File:Occupancy_9x9.png]]|}{| class="wikitable mw-collapsible mw-collapsed"! For 16x16:|-|[[File:Occupancy_16x16.png]]|}{| class="wikitable mw-collapsible mw-collapsed"! For 25x25:|-|[[File:Occupancy_25x25.png]]|}
