Заготовка wave_step_shared

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++;