Параллельность на уровне Cuda

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 d = blockIdx.x * blockDim.x + threadIdx.x;
-        int offset = day_offsets[d];
-        int count = day_counts[d];
     
-        GpuDayStats stats;
+    if (d >= num_days) return;
-        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++) {
+    int offset = day_offsets[d];
-            const GpuRecord& r = records[offset + i];
+    int count = day_counts[d];
     
-            // min/max
+    GpuDayStats stats;
-            if (r.low < stats.low) stats.low = r.low;
+    stats.day = day_indices[d];
-            if (r.high > stats.high) stats.high = r.high;
+    stats.low = DBL_MAX;
+    stats.high = -DBL_MAX;
+    stats.first_ts = DBL_MAX;
+    stats.last_ts = -DBL_MAX;
+    stats.open = 0;
+    stats.close = 0;
     
-            // first/last по timestamp
+    for (int i = 0; i < count; i++) {
-            if (r.timestamp < stats.first_ts) {
+        const GpuRecord& r = records[offset + i];
-                stats.first_ts = r.timestamp;
+        
-                stats.open = r.open;
+        // min/max
-            }
+        if (r.low < stats.low) stats.low = r.low;
-            if (r.timestamp > stats.last_ts) {
+        if (r.high > stats.high) stats.high = r.high;
-                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 поток)
+    // Запускаем kernel: каждый поток обрабатывает один день
-    aggregate_kernel<<<1, 1>>>(
+    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