From 3c707eabb3de306c11f64c5220f336bd5535be1d Mon Sep 17 00:00:00 2001
From: Arity-T <tish.mainmail@gmail.com>
Date: Sun, 30 Mar 2025 21:06:55 +0300
Subject: [PATCH] =?UTF-8?q?=D0=97=D0=B0=D0=B3=D0=BE=D1=82=D0=BE=D0=B2?=
 =?UTF-8?q?=D0=BA=D0=B0=20wave=5Fstep=5Fshared?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

---
 kernel.cu | 87 +++++++++++++++++++++++++++++++++++++++++++++++++++----
 1 file changed, 82 insertions(+), 5 deletions(-)

diff --git a/kernel.cu b/kernel.cu
index 5f4809d..94d86c7 100644
--- a/kernel.cu
+++ b/kernel.cu
@@ -4,17 +4,22 @@
 #include <cuda_runtime.h>
 #include <device_launch_parameters.h>
 
+// Настройки эксперимента
+#define USE_SHARED_MEMORY false
 #define BLOCKS_COUNT 16
-#define THREADS_COUNT 16
+#define THREADS_COUNT 16 // Используется, если USE_SHARED_MEMORY == false
+#define BLOCK_SIZE 4 // Используется, если USE_SHARED_MEMORY == true
 #define MATRIX_SIZE 32
-#define OBSTACLE_PROB 10
-#define START_X 2
+
+#define OBSTACLE_PROB 10 // Процент препятствий на полигоне
+#define START_X 2 // Начальная точка
 #define START_Y 2
-#define FINISH_X (MATRIX_SIZE - 3)
+#define FINISH_X (MATRIX_SIZE - 3) // Конечная точка
 #define FINISH_Y (MATRIX_SIZE - 3)
 
 #define INF UINT_MAX // Используем беззнаковый максимум
 
+// Вспомогательные функции
 void generate_polygon(int* P, int n) {
     srand(42);
     for (int i = 0; i < n*n; i++) {
@@ -43,6 +48,7 @@ void print_distance_map(int* P, unsigned int* dist, int n) {
     }
 }
 
+// Ядро, не использующее разделяемую память
 __global__ void wave_step(int* P, unsigned int* dist, int n, bool* changed) {
     int tid = threadIdx.x + blockIdx.x * blockDim.x;
     // printf("Hello from CUDA kernel! I'm thread #%d\n", tid);
@@ -82,6 +88,72 @@ __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) {
+    int local_row = threadIdx.x;
+    int local_col = threadIdx.y;
+
+    __shared__ bool local_changed;
+    __shared__ int shared_block[BLOCK_SIZE][BLOCK_SIZE];
+
+    int block_index = blockIdx.x;
+    const int matrix_blocks_size = (MATRIX_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE;
+
+    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;
+
+        // Заполняем shared_block в разделяемой памяти
+        copy_block_to_shared(P, dist, block_row, block_col, shared_block);
+
+        // Тот же волновой алгоритм, но в масштабах одного блока
+        do {
+            // Обновляем флаг перед началом новой итерации
+            if (threadIdx.x == 0 && threadIdx.y == 0) {
+                local_changed = false;
+            }
+
+            // Избавляемся от race conditions 
+            // (если нулевой поток отстанет, то он может невовремя сбросить флаг)
+            __syncthreads();
+
+            if (shared_block[local_row][local_col] != -1) {
+                unsigned int current_dist = dist[tid];
+                unsigned int min_dist = current_dist;
+
+                // Обновляем ячейку. Берём соседей из shared memory (shared_block), 
+                // если они там есть, иначе берём из global memory (dist)
+                
+                // ...
+            }
+
+            // Ждём, пока все потоки попробуют обновить ячейку
+            // и флаг будет в актуальном состоянии
+            __syncthreads();
+        } while (condition);
+        
+        // Каждый блок потоков обрабатывает каждый BLOCKS_COUNT блок полигона
+        block_index += BLOCKS_COUNT;
+    }
+}
+
 int main() {
     const int n = MATRIX_SIZE;
 
@@ -119,7 +191,12 @@ int main() {
         // printf("Wave step #%d\n", iterations);
         changed = false;
         cudaMemcpy(d_changed, &changed, sizeof(bool), cudaMemcpyHostToDevice);
-        wave_step<<<BLOCKS_COUNT, THREADS_COUNT>>>(d_P, d_dist, n, d_changed);
+        if (USE_SHARED_MEMORY) {
+            dim3 block(BLOCK_SIZE, BLOCK_SIZE);
+            wave_step_shared<<<BLOCKS_COUNT, block>>>(d_P, d_dist, n, d_changed);
+        } else {
+            wave_step<<<BLOCKS_COUNT, THREADS_COUNT>>>(d_P, d_dist, n, d_changed);
+        }
         cudaDeviceSynchronize(); // Синхронизация после ядра
         cudaMemcpy(&changed, d_changed, sizeof(bool), cudaMemcpyDeviceToHost);
         iterations++;