UCP/PERF: CR fixes

Author: iyastreb

src/tools/perf/cuda/cuda_kernel.cuh

@@ -88,16 +88,11 @@ __host__ UCS_F_DEVICE unsigned ucx_perf_cuda_thread_index(size_t tid)
 {
     switch (level) {
     case UCS_DEVICE_LEVEL_THREAD: return tid;
-    /* TODO: use UCS_DEVICE_NUM_THREADS_IN_WARP */
-    case UCS_DEVICE_LEVEL_WARP:   return tid / 32;
+    case UCS_DEVICE_LEVEL_WARP:   return tid / UCS_DEVICE_NUM_THREADS_IN_WARP;
     default:                      return 0;
     }
 }
 
-__host__ UCS_F_DEVICE unsigned ucx_ceil_div(unsigned x, unsigned y) {
-    return (x / y) + ((x % y) != 0);
-}
-
 #define UCX_PERF_THREAD_INDEX_SET(_level, _tid, _outval) \
     (_outval) = ucx_perf_cuda_thread_index<_level>(_tid)
 
@@ -136,8 +131,8 @@ __host__ UCS_F_DEVICE unsigned ucx_ceil_div(unsigned x, unsigned y) {
     do { \
         unsigned _blocks     = _perf.params.device_block_count; \
         unsigned _threads    = _perf.params.device_thread_count; \
-        unsigned _reqs_count = ucx_ceil_div(_perf.params.max_outstanding, \
-                                            _perf.params.device_fc_window); \
+        unsigned _reqs_count = ucs_div_round_up(_perf.params.max_outstanding, \
+                                                _perf.params.device_fc_window); \
         size_t _shared_size  = _reqs_count * sizeof(ucp_device_request_t) * \
                                ucx_perf_cuda_thread_index<_level>(_threads); \
         _kernel<_level, _cmd><<<_blocks, _threads, _shared_size>>>(__VA_ARGS__); \

src/tools/perf/cuda/ucp_cuda_kernel.cu

@@ -26,7 +26,7 @@ public:
                                   ucp_device_request_t *requests)
         : m_size(size),
           m_fc_window(fc_window),
-          m_reqs_count(ucx_ceil_div(size, fc_window)),
+          m_reqs_count(ucs_div_round_up(size, fc_window)),
           m_pending_count(0),
           m_requests(requests),
           m_pending_map(0)
@@ -357,8 +357,8 @@ ucp_perf_cuda_put_bw_kernel(ucx_perf_cuda_context &ctx,
 {
     extern __shared__ ucp_device_request_t shared_requests[];
     unsigned thread_index      = ucx_perf_cuda_thread_index<level>(threadIdx.x);
-    unsigned reqs_count        = ucx_ceil_div(ctx.max_outstanding,
-                                              ctx.device_fc_window);
+    unsigned reqs_count        = ucs_div_round_up(ctx.max_outstanding,
+                                                  ctx.device_fc_window);
     ucp_device_request_t *reqs = &shared_requests[reqs_count * thread_index];
 
     ucp_perf_cuda_request_manager req_mgr(ctx.max_outstanding,