diff --git a/kernel.cu b/kernel.cu index 94d86c7..0d9cbe6 100644 --- a/kernel.cu +++ b/kernel.cu @@ -5,7 +5,7 @@ #include // Настройки эксперимента -#define USE_SHARED_MEMORY false +#define USE_SHARED_MEMORY true #define BLOCKS_COUNT 16 #define THREADS_COUNT 16 // Используется, если USE_SHARED_MEMORY == false #define BLOCK_SIZE 4 // Используется, если USE_SHARED_MEMORY == true @@ -88,30 +88,16 @@ __global__ void wave_step(int* P, unsigned int* dist, int n, bool* changed) { } } -__device__ void copy_block_to_shared( - int* P, unsigned int* dist, int block_row, int block_col, int[][BLOCK_SIZE] shared_block -) { - int local_row = threadIdx.x; - int local_col = threadIdx.y; - int row = local_row + block_row * BLOCK_SIZE; - int col = local_col + block_col * BLOCK_SIZE; - - if (row < MATRIX_SIZE && col < MATRIX_SIZE && P[row + col * MATRIX_SIZE] != -1) { - shared_block[local_row][local_col] = dist[global_y * matrix_size + global_x]; - } else { - shared_block[local_row][local_col] = -1; - } - - __syncthreads(); // Ждём, пока все потоки скопируют данные -} - // Ядро, использующее разделяемую память -__global__ void wave_step_shared(int* P, unsigned int* dist, int n, bool* changed) { +__global__ void wave_step_shared(int* P, unsigned int* dist, int n, bool* global_changed) { int local_row = threadIdx.x; int local_col = threadIdx.y; __shared__ bool local_changed; - __shared__ int shared_block[BLOCK_SIZE][BLOCK_SIZE]; + __shared__ int shared_p[BLOCK_SIZE][BLOCK_SIZE]; + __shared__ int shared_dist[BLOCK_SIZE][BLOCK_SIZE]; + + bool changed_at_least_once = false; int block_index = blockIdx.x; const int matrix_blocks_size = (MATRIX_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE; @@ -119,9 +105,19 @@ __global__ void wave_step_shared(int* P, unsigned int* dist, int n, bool* change while (block_index < matrix_blocks_size * matrix_blocks_size) { int block_row = block_index / matrix_blocks_size; int block_col = block_index % matrix_blocks_size; + int i = local_row + block_row * BLOCK_SIZE; + int j = local_col + block_col * BLOCK_SIZE; - // Заполняем shared_block в разделяемой памяти - copy_block_to_shared(P, dist, block_row, block_col, shared_block); + // Заполняем shared_p и shared_dist в разделяемой памяти + shared_p[local_row][local_col] = 0; + shared_dist[local_row][local_col] = INF; + if (i < MATRIX_SIZE && j < MATRIX_SIZE && P[i * MATRIX_SIZE + j] != -1) { + shared_dist[local_row][local_col] = dist[i * MATRIX_SIZE + j]; + } else { + shared_p[local_row][local_col] = -1; + } + + __syncthreads(); // Ждём, пока все потоки скопируют данные // Тот же волновой алгоритм, но в масштабах одного блока do { @@ -134,24 +130,64 @@ __global__ void wave_step_shared(int* P, unsigned int* dist, int n, bool* change // (если нулевой поток отстанет, то он может невовремя сбросить флаг) __syncthreads(); - if (shared_block[local_row][local_col] != -1) { - unsigned int current_dist = dist[tid]; + if (shared_p[local_row][local_col] != -1) { + unsigned int current_dist = shared_dist[local_row][local_col]; unsigned int min_dist = current_dist; - // Обновляем ячейку. Берём соседей из shared memory (shared_block), + // Берём соседей из shared memory (shared_dist), // если они там есть, иначе берём из global memory (dist) + + // Сверху + if (local_row > 0 && shared_dist[local_row - 1][local_col] != INF) + min_dist = min(min_dist, shared_dist[local_row - 1][local_col] + 1); + else if (i > 0 && dist[(i-1)*n + j] != INF) + min_dist = min(min_dist, dist[(i-1)*n + j] + 1); + + // Снизу + if (local_row < BLOCK_SIZE - 1 && shared_dist[local_row + 1][local_col] != INF) + min_dist = min(min_dist, shared_dist[local_row + 1][local_col] + 1); + else if (i < n-1 && dist[(i+1)*n + j] != INF) + min_dist = min(min_dist, dist[(i+1)*n + j] + 1); - // ... + // Слева + if (local_col > 0 && shared_dist[local_row][local_col - 1] != INF) + min_dist = min(min_dist, shared_dist[local_row][local_col - 1] + 1); + else if (j > 0 && dist[i*n + (j-1)] != INF) + min_dist = min(min_dist, dist[i*n + (j-1)] + 1); + + // Справа + if (local_col < BLOCK_SIZE - 1 && shared_dist[local_row][local_col + 1] != INF) + min_dist = min(min_dist, shared_dist[local_row][local_col + 1] + 1); + else if (j < n-1 && dist[i*n + (j+1)] != INF) + min_dist = min(min_dist, dist[i*n + (j+1)] + 1); + + // Обновляем дистанцию в разделяемой памяти + if (min_dist < current_dist) { + shared_dist[local_row][local_col] = min_dist; + local_changed = true; + changed_at_least_once = true; + } } // Ждём, пока все потоки попробуют обновить ячейку // и флаг будет в актуальном состоянии __syncthreads(); - } while (condition); + } while (local_changed); + + // Копируем результат из разделяемой памяти (shared_dist) в глобальную (dist) + if (i < MATRIX_SIZE && j < MATRIX_SIZE && shared_p[local_row][local_col] != -1) { + dist[i * MATRIX_SIZE + j] = shared_dist[local_row][local_col]; + } // Каждый блок потоков обрабатывает каждый BLOCKS_COUNT блок полигона block_index += BLOCKS_COUNT; } + + // Если блок потоков обновил хотя бы одно значение, то, возможно, + // нужны ещё глобальные итерации + if (changed_at_least_once) { + *global_changed = true; + } } int main() { @@ -218,8 +254,14 @@ int main() { // Вывод результатов printf("Time: %.2f ms\n", milliseconds); - printf("Matrix: %dx%d | BlocksXThreads: %dx%d | Obstacles: %d%%\n\n", + if (USE_SHARED_MEMORY) { + printf("Results of the algorithm using shared memory.\n"); + printf("Matrix: %dx%d | BlocksXThreadXThread: %dx%dx%d | Obstacles: %d%%\n\n", + n, n, BLOCKS_COUNT, BLOCK_SIZE, BLOCK_SIZE, OBSTACLE_PROB); + } else { + printf("Matrix: %dx%d | BlocksXThreads: %dx%d | Obstacles: %d%%\n\n", n, n, BLOCKS_COUNT, THREADS_COUNT, OBSTACLE_PROB); + } if (MATRIX_SIZE <= 100) print_distance_map(P, dist, MATRIX_SIZE);