TOOLS/PERF: Perftest cuda kernel fixes

Author: yosefe

src/tools/perf/cuda/cuda_kernel.cuh

@@ -47,13 +47,9 @@ ucx_perf_cuda_update_report(ucx_perf_cuda_context &ctx,
     }
 }
 
-template <typename Base>
-class ucx_perf_cuda_test_runner: public Base {
+class ucx_perf_cuda_test_runner {
 public:
-    using psn_t = uint64_t;
-    using Base::m_perf;
-
-    ucx_perf_cuda_test_runner(ucx_perf_context_t &perf) : Base(perf)
+    ucx_perf_cuda_test_runner(ucx_perf_context_t &perf) : m_perf(perf)
     {
         ucs_status_t status = init_ctx();
         if (status != UCS_OK) {
@@ -62,12 +58,15 @@ public:
         }
 
         m_cpu_ctx->max_outstanding    = perf.params.max_outstanding;
-        m_cpu_ctx->max_iters          = perf.params.max_iter;
-        m_cpu_ctx->report_interval_ns = perf.params.report_interval *
-                                        UCS_NSEC_PER_SEC;
+        m_cpu_ctx->max_iters          = perf.max_iter;
         m_cpu_ctx->completed_iters    = 0;
-
-        m_poll_interval               = perf.params.report_interval / 10000;
+        if (perf.report_interval == ULONG_MAX) {
+            m_cpu_ctx->report_interval_ns = ULONG_MAX;
+        } else {
+            m_cpu_ctx->report_interval_ns = ucs_time_to_nsec(
+                                                    perf.report_interval) /
+                                            100;
+        }
     }
 
     ~ucx_perf_cuda_test_runner()
@@ -77,54 +76,26 @@ public:
 
     ucx_perf_cuda_context &gpu_ctx() const { return *m_gpu_ctx; }
 
-    UCS_F_ALWAYS_INLINE psn_t get_sn(const psn_t *gpu_ptr, const psn_t *cpu_ptr)
-    {
-        if (cpu_ptr != nullptr) {
-            return *cpu_ptr;
-        }
-
-        unsigned my_index          = rte_call(&m_perf, group_index);
-        ucs_memory_type_t mem_type = my_index ? m_perf.params.send_mem_type :
-                                                m_perf.params.recv_mem_type;
-        auto allocator             = my_index ? m_perf.send_allocator :
-                                                m_perf.recv_allocator;
-        return Base::get_sn(gpu_ptr, mem_type, allocator);
-    }
-
-    psn_t wait_sn_geq(const psn_t *gpu_ptr, const psn_t *cpu_ptr, psn_t value)
+    void wait_for_kernel(size_t msg_length)
     {
-        psn_t sn = get_sn(gpu_ptr, cpu_ptr);
-        if (sn >= value) {
-            return sn;
-        }
-
-        // TODO: use cuStreamWaitValue64 if available
-        usleep(m_poll_interval);
-        return get_sn(gpu_ptr, cpu_ptr);
-    }
-
-    void wait_for_kernel(size_t length)
-    {
-        psn_t last_completed = 0;
-        while (last_completed < m_perf.params.max_iter) {
-            psn_t completed = wait_sn_geq(&m_gpu_ctx->completed_iters,
-                                          &m_cpu_ctx->completed_iters,
-                                          last_completed);
-            psn_t delta     = completed - last_completed;
+        ucx_perf_counter_t last_completed = 0;
+        ucx_perf_counter_t completed      = m_cpu_ctx->completed_iters;
+        while (1) {
+            ucx_perf_counter_t delta = completed - last_completed;
             if (delta > 0) {
                 // TODO: calculate latency percentile on kernel
-                ucx_perf_update_multi(&m_perf, delta, delta * length);
+                ucx_perf_update(&m_perf, delta, msg_length);
+            } else if (completed >= m_perf.max_iter) {
+                break;
             }
             last_completed = completed;
+            completed      = m_cpu_ctx->completed_iters;
+            usleep(100);
         }
     }
 
-    void wait_for_sn(size_t length)
-    {
-        const psn_t *ptr = Base::sn_ptr(m_perf.recv_buffer, length);
-        while (wait_sn_geq(ptr, nullptr, m_perf.params.max_iter)
-               < m_perf.params.max_iter);
-    }
+protected:
+    ucx_perf_context_t &m_perf;
 
 private:
     ucs_status_t init_ctx()
@@ -150,7 +121,6 @@ private:
 
     ucx_perf_cuda_context *m_cpu_ctx;
     ucx_perf_cuda_context *m_gpu_ctx;
-    double                m_poll_interval;
 };
 
 

src/tools/perf/cuda/ucp_cuda_kernel.cu

@@ -23,7 +23,7 @@ ucp_perf_cuda_put_multi_bw_kernel(ucx_perf_cuda_context &ctx)
 
     for (ucx_perf_counter_t idx = 0; idx < max_iters; idx++) {
         // TODO: replace with actual put multi call
-        __nanosleep(1000000); // 1ms
+        __nanosleep(100000); // 100us
 
         ucx_perf_cuda_update_report(ctx, idx + 1, max_iters, last_report_time);
         __syncthreads();
@@ -40,23 +40,19 @@ ucp_perf_cuda_put_multi_latency_kernel(ucx_perf_cuda_context &ctx, bool is_sende
     for (ucx_perf_counter_t idx = 0; idx < max_iters; idx++) {
         // TODO: replace with actual put multi call
         // TODO: wait for completion
-        __nanosleep(1000000); // 1ms
+        __nanosleep(100000); // 100us
 
         ucx_perf_cuda_update_report(ctx, idx + 1, max_iters, last_report_time);
         __syncthreads();
     }
 }
 
-class ucp_perf_cuda_test_runner:
-    public ucx_perf_cuda_test_runner<ucp_perf_test_runner_base<uint64_t>> {
+class ucp_perf_cuda_test_runner : public ucx_perf_cuda_test_runner {
 public:
-    using psn_t = uint64_t;
-
     ucp_perf_cuda_test_runner(ucx_perf_context_t &perf) :
-        ucx_perf_cuda_test_runner<ucp_perf_test_runner_base<uint64_t>>(perf)
+        ucx_perf_cuda_test_runner(perf)
     {
         size_t length = ucx_perf_get_message_size(&m_perf.params);
-        ucs_assert(length >= sizeof(psn_t));
 
         m_perf.send_allocator->memset(m_perf.send_buffer, 0, length);
         m_perf.recv_allocator->memset(m_perf.recv_buffer, 0, length);
@@ -74,8 +70,9 @@ public:
         ucp_perf_cuda_put_multi_latency_kernel
             <UCP_DEVICE_LEVEL_BLOCK><<<1, thread_count>>>(gpu_ctx(), my_index);
         CUDA_CALL(UCS_ERR_NO_DEVICE, cudaGetLastError);
-
         wait_for_kernel(length);
+
+        CUDA_CALL(UCS_ERR_IO_ERROR, cudaDeviceSynchronize);
         ucx_perf_get_time(&m_perf);
         ucp_perf_barrier(&m_perf);
         return UCS_OK;
@@ -94,36 +91,15 @@ public:
             ucp_perf_cuda_put_multi_bw_kernel<UCP_DEVICE_LEVEL_BLOCK>
                                              <<<1, thread_count>>>(gpu_ctx());
             CUDA_CALL(UCS_ERR_NO_DEVICE, cudaGetLastError);
-
             wait_for_kernel(length);
-
-            // TODO: remove once real GDAKI is used
-            send_signal(length);
-        } else if (my_index == 0) {
-            wait_for_sn(length);
         }
+        // TODO run receiver kernel
 
+        CUDA_CALL(UCS_ERR_IO_ERROR, cudaDeviceSynchronize);
         ucx_perf_get_time(&m_perf);
         ucp_perf_barrier(&m_perf);
         return UCS_OK;
     }
-
-private:
-    // TODO: remove once real GDAKI is used
-    void send_signal(size_t length)
-    {
-        ucs_memory_type_t mem_type = m_perf.params.send_mem_type;
-        write_sn(m_perf.send_buffer, mem_type, length, m_perf.params.max_iter,
-                 m_perf.ucp.self_send_rkey);
-
-        ucs_status_ptr_t request;
-        ucp_request_param_t param = {0};
-        request = ucp_put_nbx(m_perf.ucp.ep, m_perf.send_buffer, length,
-                              m_perf.ucp.remote_addr, m_perf.ucp.rkey, &param);
-        request_wait(request, mem_type, "write_sn");
-        request = ucp_ep_flush_nbx(m_perf.ucp.self_ep, &param);
-        request_wait(request, mem_type, "flush write_sn");
-    }
 };
 
 ucx_perf_device_dispatcher_t ucx_perf_cuda_dispatcher;

src/tools/perf/lib/libperf_int.h

@@ -228,24 +228,27 @@ static inline void ucx_perf_omp_barrier(ucx_perf_context_t *perf)
 
 static UCS_F_ALWAYS_INLINE void ucx_perf_update(ucx_perf_context_t *perf,
                                                 ucx_perf_counter_t iters,
-                                                size_t bytes)
+                                                size_t bytes_per_iter)
 {
     perf->current.time   = ucs_get_time();
     perf->current.iters += iters;
-    perf->current.bytes += bytes;
-    perf->current.msgs  += 1;
+    perf->current.bytes += bytes_per_iter * iters;
+    perf->current.msgs  += iters;
 
-    perf->timing_queue[perf->timing_queue_head] =
-                    perf->current.time - perf->prev_time;
-    ++perf->timing_queue_head;
-    if (perf->timing_queue_head == TIMING_QUEUE_SIZE) {
-        perf->timing_queue_head = 0;
+    if (iters == 1) {
+        perf->timing_queue[perf->timing_queue_head] = perf->current.time -
+                                                      perf->prev_time;
+        ++perf->timing_queue_head;
+        if (perf->timing_queue_head == TIMING_QUEUE_SIZE) {
+            perf->timing_queue_head = 0;
+        }
     }
 
     perf->prev_time = perf->current.time;
 
     if (ucs_unlikely((perf->current.time - perf->prev.time) >=
-                     perf->report_interval)) {
+                     perf->report_interval) &&
+        (perf->current.iters < perf->max_iter)) {
         ucx_perf_report(perf);
     }
 }

src/tools/perf/perftest_params.c

@@ -637,6 +637,12 @@ ucs_status_t adjust_test_params(perftest_params_t *params,
         params->super.max_outstanding = test->window_size;
     }
 
+    if (params->super.send_device.mem_type != UCS_MEMORY_TYPE_LAST) {
+        /* TODO: Add getter function for thread count */
+        params->super.device_thread_count = params->super.thread_count;
+        params->super.thread_count        = 1;
+    }
+
     return UCS_OK;
 }
 
@@ -847,12 +853,6 @@ ucs_status_t parse_opts(struct perftest_context *ctx, int mpi_initialized,
         }
     }
 
-    if (ctx->params.super.send_device.mem_type != UCS_MEMORY_TYPE_LAST) {
-        /* TODO: Add getter function for thread count */
-        ctx->params.super.device_thread_count = ctx->params.super.thread_count;
-        ctx->params.super.thread_count        = 1;
-    }
-
     return init_daemon_params(&ctx->params.super);
 
 err: