Changes

Jump to: navigation, search

TriForce

1,912 bytes added, 13:17, 8 April 2019
Kernel Optimization Attempts
[[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 ====
|}
Reduced superSolve runtime from 5.2 to 3.8ms
[[File:Unoptimized_vs_Optimized.png]]
Change : Replaces the boolean array hasSeen with a single int & uses bitwise operators
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 (Kernel)
|-
|
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 (Kernel)
|-
|
}
|}
 
Change : Quickly finds one section that requires a single value in one spot, by checking all sections at once and remembering a single section
Theory : Similar to the previous Kernel, trying to remove the second loop
Result : Surprisingly slow, gains little benefit from the section logic and shared memory, yet is still required to count all values
{| class="wikitable mw-collapsible mw-collapsed"
! Notify (Kernel) - Determines a single section that has a limited value (removes section loop)
|-
|
}
|}
 
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 - Counts down instead of revisiting numbersusing Int as Boolean Array(EDITED now 4.28 seconds)
|-
|
int notSeen = 0; //Boolean Array as an Integer
if (gridIdx == 0)
changed = true;
rowHas[col][row] = false;
colHas[col][row] = false;
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 section
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;
}
|}
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; bool inSection = true; int b_shuttle = 1; for (int idx = 0; idx < N; ++idx) { // Ensures that every square in each section is working on Change : uses countdown logic with 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; --count; rowCount[row][num]--; colCount[col][num]--; boxCount[box][num]--; } else if (inSection) { guess = num; } } __syncthreads(); if ((b_shuttle & notSeen) && (rowCount[row][num] == 1 || boxCount[box][num] == 1 || colCount[col][num] == 1)) inSection = false;  b_shuttle <<= 1; }  if (count == 1 || !inSection) { at = guess + 1; notSeen = count = 0; rowHas[row][guess] = true; colHas[col][guess] = true; boxHas[box][guess] = true; changed = true; } __syncthreads();boolean array }; Result : Similar times to other Countdown kernel
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;
}
|}
{| class="wikitable mw-collapsible mw-collapsed"
! Countdown Boolean Array(EDITED - now 4.37ms)
|-
|
__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 i idx = 0; i idx < N; ++iidx) notSeen[iidx] = false; 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; 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 (changed) { __syncthreads(); if (gridIdx == 0) // forget previous change changed = false; bool inSection = true; 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 (at != UNASSIGNED || rowHas[row][num] || boxHas[box][num] || colHas[col][num])) { notSeen[num] = false; --count; rowCount[row][num]--; colCount[col][num]--; boxCount[box][num]--; } else if (inSection) { guess notSeen[num] = numfalse; } --count; } __syncthreads(); if (at == UNASSIGNED && notSeen[num] && (count == 1 || rowCount[row][num] == 1 || boxCount[box][num] == 1 || colCount[col][num] == 1)) inSection = false;  }  if (count == 1 || !inSection) { at = guess + 1; count = 0; rowHas[row][guessnum] = true; colHas[col][guessnum] = true; boxHas[box][guessnum] = 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]]

Navigation menu