From d5ca93c188e30e5e676546a34ae5b1f1319a9f78 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 22 Jan 2025 11:31:36 +0100 Subject: [PATCH] Backport to 2.8: Deprecate cub::DeviceSpmv (#3320) (#3374) Fixes: #896 --- cub/cub/agent/agent_spmv_orig.cuh | 15 +++++++-- cub/cub/device/device_spmv.cuh | 31 +++++++++-------- .../device/dispatch/dispatch_spmv_orig.cuh | 33 +++++++++++++++++-- cub/test/test_device_spmv.cu | 4 +++ 4 files changed, 64 insertions(+), 19 deletions(-) diff --git a/cub/cub/agent/agent_spmv_orig.cuh b/cub/cub/agent/agent_spmv_orig.cuh index 2ad0bee84a6..80d571d58db 100644 --- a/cub/cub/agent/agent_spmv_orig.cuh +++ b/cub/cub/agent/agent_spmv_orig.cuh @@ -104,7 +104,7 @@ template -struct AgentSpmvPolicy +struct CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") AgentSpmvPolicy { enum { @@ -150,7 +150,12 @@ struct AgentSpmvPolicy * Signed integer type for sequence offsets */ template -struct SpmvParams +struct +// with NVHPC, we get a deprecation warning in the implementation of cudaLaunchKernelEx, which we cannot suppress :/ +#if !_CCCL_COMPILER(NVHPC) + CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") +#endif + SpmvParams { /// Pointer to the array of \p num_nonzeros values of the corresponding nonzero elements of matrix /// A. @@ -213,7 +218,7 @@ template -struct AgentSpmv +struct CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") AgentSpmv { //--------------------------------------------------------------------- // Types and constants @@ -310,7 +315,9 @@ struct AgentSpmv /// Reference to temp_storage _TempStorage& temp_storage; + _CCCL_SUPPRESS_DEPRECATED_PUSH SpmvParams& spmv_params; + _CCCL_SUPPRESS_DEPRECATED_POP /// Wrapped pointer to the array of \p num_nonzeros values of the corresponding nonzero elements /// of matrix A. @@ -343,6 +350,7 @@ struct AgentSpmv * @param spmv_params * SpMV input parameter bundle */ + _CCCL_SUPPRESS_DEPRECATED_PUSH _CCCL_DEVICE _CCCL_FORCEINLINE AgentSpmv(TempStorage& temp_storage, SpmvParams& spmv_params) : temp_storage(temp_storage.Alias()) , spmv_params(spmv_params) @@ -352,6 +360,7 @@ struct AgentSpmv , wd_vector_x(spmv_params.d_vector_x) , wd_vector_y(spmv_params.d_vector_y) {} + _CCCL_SUPPRESS_DEPRECATED_POP /** * @brief Consume a merge tile, specialized for direct-load of nonzeros diff --git a/cub/cub/device/device_spmv.cuh b/cub/cub/device/device_spmv.cuh index 114454a791a..6f8dda0c48c 100644 --- a/cub/cub/device/device_spmv.cuh +++ b/cub/cub/device/device_spmv.cuh @@ -78,7 +78,7 @@ CUB_NAMESPACE_BEGIN //! @cdp_class{DeviceSpmv} //! //! @endrst -struct DeviceSpmv +struct CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") DeviceSpmv { //! @name CSR matrix operations //! @{ @@ -177,22 +177,25 @@ struct DeviceSpmv //! **[optional]** CUDA stream to launch kernels within. Default is stream\ :sub:`0`. //! @endrst template - CUB_RUNTIME_FUNCTION static cudaError_t CsrMV( - void* d_temp_storage, - size_t& temp_storage_bytes, - const ValueT* d_values, - const int* d_row_offsets, - const int* d_column_indices, - const ValueT* d_vector_x, - ValueT* d_vector_y, - int num_rows, - int num_cols, - int num_nonzeros, - cudaStream_t stream = 0) + CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") + CUB_RUNTIME_FUNCTION static cudaError_t + CsrMV(void* d_temp_storage, + size_t& temp_storage_bytes, + const ValueT* d_values, + const int* d_row_offsets, + const int* d_column_indices, + const ValueT* d_vector_x, + ValueT* d_vector_y, + int num_rows, + int num_cols, + int num_nonzeros, + cudaStream_t stream = 0) { CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceSpmv::CsrMV"); + _CCCL_SUPPRESS_DEPRECATED_PUSH SpmvParams spmv_params; + _CCCL_SUPPRESS_DEPRECATED_POP spmv_params.d_values = d_values; spmv_params.d_row_end_offsets = d_row_offsets + 1; spmv_params.d_column_indices = d_column_indices; @@ -204,7 +207,9 @@ struct DeviceSpmv spmv_params.alpha = ValueT{1}; spmv_params.beta = ValueT{0}; + _CCCL_SUPPRESS_DEPRECATED_PUSH return DispatchSpmv::Dispatch(d_temp_storage, temp_storage_bytes, spmv_params, stream); + _CCCL_SUPPRESS_DEPRECATED_POP } #ifndef _CCCL_DOXYGEN_INVOKED // Do not document diff --git a/cub/cub/device/dispatch/dispatch_spmv_orig.cuh b/cub/cub/device/dispatch/dispatch_spmv_orig.cuh index b999b109d2e..61df040051b 100644 --- a/cub/cub/device/dispatch/dispatch_spmv_orig.cuh +++ b/cub/cub/device/dispatch/dispatch_spmv_orig.cuh @@ -62,6 +62,7 @@ #include +_CCCL_SUPPRESS_DEPRECATED_PUSH CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -83,7 +84,9 @@ CUB_NAMESPACE_BEGIN * @param[in] spmv_params * SpMV input parameter bundle */ +_CCCL_SUPPRESS_DEPRECATED_PUSH template +CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmv1ColKernel(SpmvParams spmv_params) { using VectorValueIteratorT = @@ -106,6 +109,7 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmv1ColKernel(SpmvParams -CUB_DETAIL_KERNEL_ATTRIBUTES void -DeviceSpmvSearchKernel(int num_merge_tiles, CoordinateT* d_tile_coordinates, SpmvParamsT spmv_params) +CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") +CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmvSearchKernel( + int num_merge_tiles, CoordinateT* d_tile_coordinates, SpmvParamsT spmv_params) { /// Constants enum @@ -210,6 +215,7 @@ DeviceSpmvSearchKernel(int num_merge_tiles, CoordinateT* d_tile_coordinates, Spm * @param[in] num_segment_fixup_tiles * Number of reduce-by-key tiles (fixup grid size) */ +_CCCL_SUPPRESS_DEPRECATED_PUSH template +CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") __launch_bounds__(int(SpmvPolicyT::BLOCK_THREADS)) CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmvKernel( SpmvParams spmv_params, CoordinateT* d_tile_coordinates, @@ -226,7 +233,9 @@ __launch_bounds__(int(SpmvPolicyT::BLOCK_THREADS)) CUB_DETAIL_KERNEL_ATTRIBUTES int num_segment_fixup_tiles) { // Spmv agent type specialization + _CCCL_SUPPRESS_DEPRECATED_PUSH using AgentSpmvT = AgentSpmv; + _CCCL_SUPPRESS_DEPRECATED_POP // Shared memory for AgentSpmv __shared__ typename AgentSpmvT::TempStorage temp_storage; @@ -236,6 +245,7 @@ __launch_bounds__(int(SpmvPolicyT::BLOCK_THREADS)) CUB_DETAIL_KERNEL_ATTRIBUTES // Initialize fixup tile status tile_state.InitializeStatus(num_segment_fixup_tiles); } +_CCCL_SUPPRESS_DEPRECATED_POP /** * @tparam ValueT @@ -247,7 +257,9 @@ __launch_bounds__(int(SpmvPolicyT::BLOCK_THREADS)) CUB_DETAIL_KERNEL_ATTRIBUTES * @tparam HAS_BETA * Whether the input parameter Beta is 0 */ +_CCCL_SUPPRESS_DEPRECATED_PUSH template +CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmvEmptyMatrixKernel(SpmvParams spmv_params) { const int row = static_cast(threadIdx.x + blockIdx.x * blockDim.x); @@ -264,6 +276,7 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmvEmptyMatrixKernel(SpmvParams +CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") __launch_bounds__(int(AgentSegmentFixupPolicyT::BLOCK_THREADS)) CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSegmentFixupKernel( PairsInputIteratorT d_pairs_in, @@ -327,6 +342,7 @@ __launch_bounds__(int(AgentSegmentFixupPolicyT::BLOCK_THREADS)) AgentSegmentFixupT(temp_storage, d_pairs_in, d_aggregates_out, ::cuda::std::equal_to<>{}, ::cuda::std::plus<>{}) .ConsumeRange(num_items, num_tiles, tile_state); } +_CCCL_SUPPRESS_DEPRECATED_POP /****************************************************************************** * Dispatch @@ -342,7 +358,7 @@ __launch_bounds__(int(AgentSegmentFixupPolicyT::BLOCK_THREADS)) * Signed integer type for global offsets */ template -struct DispatchSpmv +struct CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") DispatchSpmv { //--------------------------------------------------------------------- // Constants and Types @@ -355,7 +371,9 @@ struct DispatchSpmv }; // SpmvParams bundle type + _CCCL_SUPPRESS_DEPRECATED_PUSH using SpmvParamsT = SpmvParams; + _CCCL_SUPPRESS_DEPRECATED_POP // 2D merge path coordinate type using CoordinateT = typename CubVector::Type; @@ -373,6 +391,7 @@ struct DispatchSpmv /// SM35 struct Policy350 { + _CCCL_SUPPRESS_DEPRECATED_PUSH using SpmvPolicyT = AgentSpmvPolicy<(sizeof(ValueT) > 4) ? 96 : 128, (sizeof(ValueT) > 4) ? 4 : 7, @@ -383,6 +402,7 @@ struct DispatchSpmv LOAD_LDG, (sizeof(ValueT) > 4) ? true : false, BLOCK_SCAN_WARP_SCANS>; + _CCCL_SUPPRESS_DEPRECATED_POP using SegmentFixupPolicyT = AgentSegmentFixupPolicy<128, 3, BLOCK_LOAD_VECTORIZE, LOAD_LDG, BLOCK_SCAN_WARP_SCANS>; }; @@ -390,6 +410,7 @@ struct DispatchSpmv /// SM37 struct Policy370 { + _CCCL_SUPPRESS_DEPRECATED_PUSH using SpmvPolicyT = AgentSpmvPolicy<(sizeof(ValueT) > 4) ? 128 : 128, (sizeof(ValueT) > 4) ? 9 : 14, @@ -400,6 +421,7 @@ struct DispatchSpmv LOAD_LDG, false, BLOCK_SCAN_WARP_SCANS>; + _CCCL_SUPPRESS_DEPRECATED_POP using SegmentFixupPolicyT = AgentSegmentFixupPolicy<128, 3, BLOCK_LOAD_VECTORIZE, LOAD_LDG, BLOCK_SCAN_WARP_SCANS>; }; @@ -407,6 +429,7 @@ struct DispatchSpmv /// SM50 struct Policy500 { + _CCCL_SUPPRESS_DEPRECATED_PUSH using SpmvPolicyT = AgentSpmvPolicy<(sizeof(ValueT) > 4) ? 64 : 128, (sizeof(ValueT) > 4) ? 6 : 7, @@ -417,6 +440,7 @@ struct DispatchSpmv LOAD_LDG, (sizeof(ValueT) > 4) ? true : false, (sizeof(ValueT) > 4) ? BLOCK_SCAN_WARP_SCANS : BLOCK_SCAN_RAKING_MEMOIZE>; + _CCCL_SUPPRESS_DEPRECATED_POP using SegmentFixupPolicyT = AgentSegmentFixupPolicy<128, 3, BLOCK_LOAD_VECTORIZE, LOAD_LDG, BLOCK_SCAN_RAKING_MEMOIZE>; @@ -425,6 +449,7 @@ struct DispatchSpmv /// SM60 struct Policy600 { + _CCCL_SUPPRESS_DEPRECATED_PUSH using SpmvPolicyT = AgentSpmvPolicy<(sizeof(ValueT) > 4) ? 64 : 128, (sizeof(ValueT) > 4) ? 5 : 7, @@ -435,6 +460,7 @@ struct DispatchSpmv LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS>; + _CCCL_SUPPRESS_DEPRECATED_POP using SegmentFixupPolicyT = AgentSegmentFixupPolicy<128, 3, BLOCK_LOAD_DIRECT, LOAD_LDG, BLOCK_SCAN_WARP_SCANS>; }; @@ -1005,4 +1031,5 @@ struct DispatchSpmv #endif // _CCCL_DOXYGEN_INVOKED }; +_CCCL_SUPPRESS_DEPRECATED_POP CUB_NAMESPACE_END diff --git a/cub/test/test_device_spmv.cu b/cub/test/test_device_spmv.cu index 5a120e56e96..0e8c932f240 100644 --- a/cub/test/test_device_spmv.cu +++ b/cub/test/test_device_spmv.cu @@ -47,6 +47,10 @@ #include #include +// note: there is no matching _CCCL_SUPPRESS_DEPRECATED_POP at the end of the file so warnings coming from +// cudafe1.stub.c file are suppressed as well +_CCCL_SUPPRESS_DEPRECATED_PUSH + bool g_verbose = false; //==============================================================================