Changes

Jump to: navigation, search

TriForce

1,136 bytes added, 13:51, 6 April 2019
Kernel Optimization Attempts
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)
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; bool inSection = true; 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; --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(); }; 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
|-
|
__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 = 0; i < N; ++i) notSeen[i] = false; if (gridIdx == 0) changed = true;
// 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 = 0; i < N; ++i) notSeen[i] = 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; 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 (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 = num; } } __syncthreads(); if (at == UNASSIGNED && notSeen[num] && (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][guess] = true; colHas[col][guess] = true; boxHas[box][guess] = true; changed = true; } __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