Deprecate cub::DeviceSpmv

Fixes: NVIDIA#896
bernhardmgruber · Jan 13, 2025 · a80c12d · a80c12d
1 parent 804dd15
commit a80c12d
Show file tree

Hide file tree

Showing 4 changed files with 36 additions and 20 deletions.
diff --git a/cub/cub/agent/agent_spmv_orig.cuh b/cub/cub/agent/agent_spmv_orig.cuh
@@ -102,7 +102,7 @@ template <int _BLOCK_THREADS,
           CacheLoadModifier _VECTOR_VALUES_LOAD_MODIFIER,
           bool _DIRECT_LOAD_NONZEROS,
           BlockScanAlgorithm _SCAN_ALGORITHM>
-struct AgentSpmvPolicy
+struct CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") AgentSpmvPolicy
 {
   enum
   {
@@ -148,7 +148,7 @@ struct AgentSpmvPolicy
  *   Signed integer type for sequence offsets
  */
 template <typename ValueT, typename OffsetT>
-struct SpmvParams
+struct CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") SpmvParams
 {
   /// Pointer to the array of \p num_nonzeros values of the corresponding nonzero elements of matrix
   /// <b>A</b>.
@@ -211,7 +211,7 @@ template <typename AgentSpmvPolicyT,
           bool HAS_ALPHA,
           bool HAS_BETA,
           int LEGACY_PTX_ARCH = 0>
-struct AgentSpmv
+struct CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") AgentSpmv
 {
   //---------------------------------------------------------------------
   // Types and constants

diff --git 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,18 +177,19 @@ struct DeviceSpmv
   //!   **[optional]** CUDA stream to launch kernels within. Default is stream\ :sub:`0`.
   //!   @endrst
   template <typename ValueT>
-  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");
 
@@ -204,7 +205,9 @@ struct DeviceSpmv
     spmv_params.alpha             = ValueT{1};
     spmv_params.beta              = ValueT{0};
 
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
     return DispatchSpmv<ValueT, int>::Dispatch(d_temp_storage, temp_storage_bytes, spmv_params, stream);
+    _CCCL_SUPPRESS_DEPRECATED_POP
   }
 
   //! @}  end member group

diff --git a/cub/cub/device/dispatch/dispatch_spmv_orig.cuh b/cub/cub/device/dispatch/dispatch_spmv_orig.cuh
@@ -84,6 +84,7 @@ CUB_NAMESPACE_BEGIN
  *   SpMV input parameter bundle
  */
 template <typename AgentSpmvPolicyT, typename ValueT, typename OffsetT>
+CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead")
 CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmv1ColKernel(SpmvParams<ValueT, OffsetT> spmv_params)
 {
   using VectorValueIteratorT =
@@ -132,8 +133,9 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmv1ColKernel(SpmvParams<ValueT, Offset
  *   SpMV input parameter bundle
  */
 template <typename SpmvPolicyT, typename OffsetT, typename CoordinateT, typename SpmvParamsT>
-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
@@ -217,6 +219,7 @@ template <typename SpmvPolicyT,
           typename CoordinateT,
           bool HAS_ALPHA,
           bool HAS_BETA>
+CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead")
 __launch_bounds__(int(SpmvPolicyT::BLOCK_THREADS)) CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmvKernel(
   SpmvParams<ValueT, OffsetT> spmv_params,
   CoordinateT* d_tile_coordinates,
@@ -226,7 +229,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<SpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA>;
+  _CCCL_SUPPRESS_DEPRECATED_POP
 
   // Shared memory for AgentSpmv
   __shared__ typename AgentSpmvT::TempStorage temp_storage;
@@ -248,6 +253,7 @@ __launch_bounds__(int(SpmvPolicyT::BLOCK_THREADS)) CUB_DETAIL_KERNEL_ATTRIBUTES
  *   Whether the input parameter Beta is 0
  */
 template <typename ValueT, typename OffsetT, bool HAS_BETA>
+CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead")
 CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmvEmptyMatrixKernel(SpmvParams<ValueT, OffsetT> spmv_params)
 {
   const int row = static_cast<int>(threadIdx.x + blockIdx.x * blockDim.x);
@@ -298,18 +304,21 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmvEmptyMatrixKernel(SpmvParams<ValueT,
  * @param[in] tile_state
  *   Tile status interface
  */
+_CCCL_SUPPRESS_DEPRECATED_PUSH
 template <typename AgentSegmentFixupPolicyT,
           typename PairsInputIteratorT,
           typename AggregatesOutputIteratorT,
           typename OffsetT,
           typename ScanTileStateT>
+CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead")
 __launch_bounds__(int(AgentSegmentFixupPolicyT::BLOCK_THREADS))
   CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSegmentFixupKernel(
     PairsInputIteratorT d_pairs_in,
     AggregatesOutputIteratorT d_aggregates_out,
     OffsetT num_items,
     int num_tiles,
-    ScanTileStateT tile_state)
+    ScanTileStateT tile_state) //
+  _CCCL_SUPPRESS_DEPRECATED_POP
 {
   // Thread block type for reducing tiles of value segments
   using AgentSegmentFixupT =
@@ -342,7 +351,7 @@ __launch_bounds__(int(AgentSegmentFixupPolicyT::BLOCK_THREADS))
  *   Signed integer type for global offsets
  */
 template <typename ValueT, typename OffsetT>
-struct DispatchSpmv
+struct CCCL_DEPRECATED_BECAUSE("Use the cuSPARSE library instead") DispatchSpmv
 {
   //---------------------------------------------------------------------
   // Constants and Types

diff --git a/cub/test/test_device_spmv.cu b/cub/test/test_device_spmv.cu
@@ -47,6 +47,8 @@
 #include <c2h/device_policy.h>
 #include <c2h/vector.h>
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+
 bool g_verbose = false;
 
 //==============================================================================
@@ -605,3 +607,5 @@ int main(int argc, char** argv)
 
   test_types();
 }
+
+_CCCL_SUPPRESS_DEPRECATED_POP