ggml-cuda: rename + remove unused var

Author: am17an

ggml/src/ggml-cuda/common.cuh

@@ -966,9 +966,8 @@ struct ggml_cuda_graph {
 };
 
 struct ggml_cuda_concurrent_event {
-    std::vector<cudaEvent_t> per_stream_events;
+    std::vector<cudaEvent_t> join_events;
     cudaEvent_t              fork_event;
-    cudaEvent_t              join_event;
 
     int                                          n_streams = 0;
     std::unordered_map<const ggml_tensor *, int> stream_mapping;
@@ -978,14 +977,13 @@ struct ggml_cuda_concurrent_event {
     ggml_cuda_concurrent_event() = default;
 
     explicit ggml_cuda_concurrent_event(int n_streams) : n_streams(n_streams) {
-        per_stream_events.resize(n_streams);
+        join_events.resize(n_streams);
 
-        for (size_t i = 0; i < per_stream_events.size(); ++i) {
-            CUDA_CHECK(cudaEventCreateWithFlags(&per_stream_events[i], cudaEventDisableTiming));
+        for (size_t i = 0; i < join_events.size(); ++i) {
+            CUDA_CHECK(cudaEventCreateWithFlags(&join_events[i], cudaEventDisableTiming));
         }
 
         CUDA_CHECK(cudaEventCreateWithFlags(&fork_event, cudaEventDisableTiming));
-        CUDA_CHECK(cudaEventCreateWithFlags(&join_event, cudaEventDisableTiming));
     }
 };
 

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -3230,9 +3230,9 @@ static void evaluate_and_capture_cuda_graph(ggml_backend_cuda_context * cuda_ctx
                         cuda_ctx->curr_stream_no = 0;
                         for (int i = 1; i <= concurrent_event->n_streams; ++i) {
                             // Wait on join events of forked streams in the main stream 
-                            CUDA_CHECK(cudaEventRecord(concurrent_event->per_stream_events[i - 1],
+                            CUDA_CHECK(cudaEventRecord(concurrent_event->join_events[i - 1],
                                                        cuda_ctx->stream(cuda_ctx->device, i)));
-                            CUDA_CHECK(cudaStreamWaitEvent(cuda_ctx->stream(), concurrent_event->per_stream_events[i - 1]));
+                            CUDA_CHECK(cudaStreamWaitEvent(cuda_ctx->stream(), concurrent_event->join_events[i - 1]));
                         }
 
                         is_concurrent_event_active = false;