From f28f81032f8a348bd908d6db1a98ebd788b485dd Mon Sep 17 00:00:00 2001 From: says1117 Date: Tue, 28 Jul 2026 03:17:01 -0400 Subject: [PATCH] Use cuda::std::numeric_limits in LAP kernels instead of passing infinity The LAP Hungarian-algorithm kernels took the sentinel infinity (std::numeric_limits::max()) as a host-computed kernel argument, guarded by FIXMEs to switch to cuda::std::numeric_limits once CUDA 10.2 was the baseline. That baseline is long met, so this computes the limit on-device instead. Drops the infinity parameter from kernel_rowReduction, kernel_columnReduction, kernel_dualUpdate_1 and kernel_dualUpdate_2 and their call sites in lap_functions.cuh, and includes cuda/std/limits. Verified: SOLVERS_TEST builds clean; all 6 Hungarian LAP tests pass. --- .../raft/solver/detail/lap_functions.cuh | 19 ++------- .../raft/solver/detail/lap_kernels.cuh | 39 ++++++------------- 2 files changed, 15 insertions(+), 43 deletions(-) diff --git a/cpp/include/raft/solver/detail/lap_functions.cuh b/cpp/include/raft/solver/detail/lap_functions.cuh index 64536ce86a..38d1d4fed0 100644 --- a/cpp/include/raft/solver/detail/lap_functions.cuh +++ b/cpp/include/raft/solver/detail/lap_functions.cuh @@ -1,6 +1,6 @@ /* * SPDX-FileCopyrightText: Copyright 2020 KETAN DATE & RAKESH NAGI - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ /* @@ -118,19 +118,14 @@ inline void initialReduction(raft::resources const& handle, detail::calculateRectangularDims(blocks_per_grid, threads_per_block, total_blocks, N, SP); kernel_rowReduction<<>>( - d_costs, d_vertices_dev.row_duals, SP, N, std::numeric_limits::max()); + d_costs, d_vertices_dev.row_duals, SP, N); RAFT_CHECK_CUDA(resource::get_cuda_stream(handle)); kernel_columnReduction<<>>( - d_costs, - d_vertices_dev.row_duals, - d_vertices_dev.col_duals, - SP, - N, - std::numeric_limits::max()); + d_costs, d_vertices_dev.row_duals, d_vertices_dev.col_duals, SP, N); RAFT_CHECK_CUDA(resource::get_cuda_stream(handle)); } @@ -495,12 +490,7 @@ inline void dualUpdate(raft::resources const& handle, detail::calculateLinearDims(blocks_per_grid, threads_per_block, total_blocks, SP); kernel_dualUpdate_1<<>>( - sp_min_v.data(), - d_vertices_dev.col_slacks, - d_vertices_dev.col_covers, - SP, - N, - std::numeric_limits::max()); + sp_min_v.data(), d_vertices_dev.col_slacks, d_vertices_dev.col_covers, SP, N); RAFT_CHECK_CUDA(resource::get_cuda_stream(handle)); @@ -516,7 +506,6 @@ inline void dualUpdate(raft::resources const& handle, d_col_data_dev.parents, SP, N, - std::numeric_limits::max(), epsilon); RAFT_CHECK_CUDA(resource::get_cuda_stream(handle)); diff --git a/cpp/include/raft/solver/detail/lap_kernels.cuh b/cpp/include/raft/solver/detail/lap_kernels.cuh index 5aa826394c..b46ca9c03f 100644 --- a/cpp/include/raft/solver/detail/lap_kernels.cuh +++ b/cpp/include/raft/solver/detail/lap_kernels.cuh @@ -1,6 +1,6 @@ /* * SPDX-FileCopyrightText: Copyright 2020 KETAN DATE & RAKESH NAGI - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ /* @@ -34,6 +34,7 @@ #include #include +#include #include #include @@ -162,15 +163,12 @@ __device__ void __augment(vertex_t* d_row_assignments, } // Kernel for reducing the rows by subtracting row minimum from each row element. -// FIXME: Once cuda 10.2 is the standard should replace passing infinity -// here with using cuda::std::numeric_limits::max() template -RAFT_KERNEL kernel_rowReduction( - weight_t const* d_costs, weight_t* d_row_duals, int SP, vertex_t N, weight_t infinity) +RAFT_KERNEL kernel_rowReduction(weight_t const* d_costs, weight_t* d_row_duals, int SP, vertex_t N) { int spid = blockIdx.y * blockDim.y + threadIdx.y; int rowid = blockIdx.x * blockDim.x + threadIdx.x; - weight_t min = infinity; + weight_t min = cuda::std::numeric_limits::max(); if (spid < SP && rowid < N) { for (int colid = 0; colid < N; colid++) { @@ -184,20 +182,14 @@ RAFT_KERNEL kernel_rowReduction( } // Kernel for reducing the column by subtracting column minimum from each column element. -// FIXME: Once cuda 10.2 is the standard should replace passing infinity -// here with using cuda::std::numeric_limits::max() template -RAFT_KERNEL kernel_columnReduction(weight_t const* d_costs, - weight_t const* d_row_duals, - weight_t* d_col_duals, - int SP, - vertex_t N, - weight_t infinity) +RAFT_KERNEL kernel_columnReduction( + weight_t const* d_costs, weight_t const* d_row_duals, weight_t* d_col_duals, int SP, vertex_t N) { int spid = blockIdx.y * blockDim.y + threadIdx.y; int colid = blockIdx.x * blockDim.x + threadIdx.x; - weight_t min = infinity; + weight_t min = cuda::std::numeric_limits::max(); if (spid < SP && colid < N) { for (int rowid = 0; rowid < N; rowid++) { @@ -443,20 +435,14 @@ RAFT_KERNEL kernel_augmentation(vertex_t* d_row_assignments, } // Kernel for updating the dual values in Step 5. -// FIXME: Once cuda 10.2 is the standard should replace passing infinity -// here with using cuda::std::numeric_limits::max() template -RAFT_KERNEL kernel_dualUpdate_1(weight_t* d_sp_min, - weight_t const* d_col_slacks, - int const* d_col_covers, - int SP, - vertex_t N, - weight_t infinity) +RAFT_KERNEL kernel_dualUpdate_1( + weight_t* d_sp_min, weight_t const* d_col_slacks, int const* d_col_covers, int SP, vertex_t N) { int spid = blockIdx.x * blockDim.x + threadIdx.x; if (spid < SP) { - weight_t min = infinity; + weight_t min = cuda::std::numeric_limits::max(); for (int colid = 0; colid < N; colid++) { int index = spid * N + colid; weight_t slack = d_col_slacks[index]; @@ -471,8 +457,6 @@ RAFT_KERNEL kernel_dualUpdate_1(weight_t* d_sp_min, } // Kernel for updating the dual values in Step 5. -// FIXME: Once cuda 10.2 is the standard should replace passing infinity -// here with using cuda::std::numeric_limits::max() template RAFT_KERNEL kernel_dualUpdate_2(weight_t const* d_sp_min, weight_t* d_row_duals, @@ -484,7 +468,6 @@ RAFT_KERNEL kernel_dualUpdate_2(weight_t const* d_sp_min, vertex_t* d_col_parents, int SP, vertex_t N, - weight_t infinity, weight_t epsilon) { int spid = blockIdx.y * blockDim.y + threadIdx.y; @@ -493,7 +476,7 @@ RAFT_KERNEL kernel_dualUpdate_2(weight_t const* d_sp_min, if (spid < SP && id < N) { int index = spid * N + id; - if (d_sp_min[spid] < infinity) { + if (d_sp_min[spid] < cuda::std::numeric_limits::max()) { weight_t theta = d_sp_min[spid]; int row_cover = d_row_covers[index]; int col_cover = d_col_covers[index];