univer-6th-semester/supercomputers
Code Actions Releases 1

shared версия алгоритма заработала (я не ожидал)

Browse Source
This commit is contained in:
Артём Тищенко
2025-03-31 11:13:03 +03:00
parent 3c707eabb3
commit a33e38a105
1 changed files with 70 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files

98
kernel.cu
View File

@@ -5,7 +5,7 @@
#include <device_launch_parameters.h>
// Настройки эксперимента
#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);