From 68ea345a35b250cceae54c9b24d9622226800d0e Mon Sep 17 00:00:00 2001
From: Arity-T <tish.mainmail@gmail.com>
Date: Tue, 2 Dec 2025 12:58:49 +0000
Subject: [PATCH] =?UTF-8?q?=D0=9F=D0=B0=D1=80=D0=B0=D0=BB=D0=BB=D0=B5?=
 =?UTF-8?q?=D0=BB=D1=8C=D0=BD=D0=BE=D1=81=D1=82=D1=8C=20=D0=BD=D0=B0=20?=
 =?UTF-8?q?=D1=83=D1=80=D0=BE=D0=B2=D0=BD=D0=B5=20Cuda?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

---
 src/gpu_plugin.cu | 73 +++++++++++++++++++++++++----------------------
 1 file changed, 39 insertions(+), 34 deletions(-)

diff --git a/src/gpu_plugin.cu b/src/gpu_plugin.cu
index 9951a0a..674775e 100644
--- a/src/gpu_plugin.cu
+++ b/src/gpu_plugin.cu
@@ -30,7 +30,7 @@ extern "C" int gpu_is_available() {
     return (n > 0) ? 1 : 0;
 }
 
-// Kernel для агрегации (один поток обрабатывает все данные)
+// Kernel для агрегации (каждый поток обрабатывает один день)
 __global__ void aggregate_kernel(
     const GpuRecord* records,
     int num_records,
@@ -40,41 +40,43 @@ __global__ void aggregate_kernel(
     int num_days,
     GpuDayStats* out_stats)
 {
-    // Один поток обрабатывает все дни последовательно
-    for (int d = 0; d < num_days; d++) {
-        int offset = day_offsets[d];
-        int count = day_counts[d];
+    // Глобальный индекс потока = индекс дня
+    int d = blockIdx.x * blockDim.x + threadIdx.x;
+    
+    if (d >= num_days) return;
+    
+    int offset = day_offsets[d];
+    int count = day_counts[d];
+    
+    GpuDayStats stats;
+    stats.day = day_indices[d];
+    stats.low = DBL_MAX;
+    stats.high = -DBL_MAX;
+    stats.first_ts = DBL_MAX;
+    stats.last_ts = -DBL_MAX;
+    stats.open = 0;
+    stats.close = 0;
+    
+    for (int i = 0; i < count; i++) {
+        const GpuRecord& r = records[offset + i];
         
-        GpuDayStats stats;
-        stats.day = day_indices[d];
-        stats.low = DBL_MAX;
-        stats.high = -DBL_MAX;
-        stats.first_ts = DBL_MAX;
-        stats.last_ts = -DBL_MAX;
-        stats.open = 0;
-        stats.close = 0;
+        // min/max
+        if (r.low < stats.low) stats.low = r.low;
+        if (r.high > stats.high) stats.high = r.high;
         
-        for (int i = 0; i < count; i++) {
-            const GpuRecord& r = records[offset + i];
-            
-            // min/max
-            if (r.low < stats.low) stats.low = r.low;
-            if (r.high > stats.high) stats.high = r.high;
-            
-            // first/last по timestamp
-            if (r.timestamp < stats.first_ts) {
-                stats.first_ts = r.timestamp;
-                stats.open = r.open;
-            }
-            if (r.timestamp > stats.last_ts) {
-                stats.last_ts = r.timestamp;
-                stats.close = r.close;
-            }
+        // first/last по timestamp
+        if (r.timestamp < stats.first_ts) {
+            stats.first_ts = r.timestamp;
+            stats.open = r.open;
+        }
+        if (r.timestamp > stats.last_ts) {
+            stats.last_ts = r.timestamp;
+            stats.close = r.close;
         }
-        
-        stats.avg = (stats.low + stats.high) / 2.0;
-        out_stats[d] = stats;
     }
+    
+    stats.avg = (stats.low + stats.high) / 2.0;
+    out_stats[d] = stats;
 }
 
 // Функция агрегации, вызываемая из C++
@@ -124,8 +126,11 @@ extern "C" int gpu_aggregate_days(
     err = cudaMemcpy(d_day_indices, h_day_indices, num_days * sizeof(long long), cudaMemcpyHostToDevice);
     if (err != cudaSuccess) return -13;
     
-    // Запускаем kernel (1 блок, 1 поток)
-    aggregate_kernel<<<1, 1>>>(
+    // Запускаем kernel: каждый поток обрабатывает один день
+    const int THREADS_PER_BLOCK = 256;
+    int num_blocks = (num_days + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
+    
+    aggregate_kernel<<<num_blocks, THREADS_PER_BLOCK>>>(
         d_records, num_records,
         d_day_offsets, d_day_counts, d_day_indices,
         num_days, d_out_stats