Skip to content

Commit

Permalink
Backport to 2.8: Deprecate cub::DeviceSpmv (#3320) (#3374)
Browse files Browse the repository at this point in the history 
Fixes: #896
  • Loading branch information
@bernhardmgruber
bernhardmgruber authored Jan 22, 2025
1 parent 67f625e commit d5ca93c
Show file tree
Hide file tree
Showing 4 changed files with 64 additions and 19 deletions.
15 changes: 12 additions & 3 deletions cub/cub/agent/agent_spmv_orig.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -104,7 +104,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
{
Expand Down Expand Up @@ -150,7 +150,12 @@ struct AgentSpmvPolicy
* Signed integer type for sequence offsets
*/
template <typename ValueT, typename OffsetT>
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
/// <b>A</b>.
Expand Down Expand Up @@ -213,7 +218,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
Expand Down Expand Up @@ -310,7 +315,9 @@ struct AgentSpmv
/// Reference to temp_storage
_TempStorage& temp_storage;

_CCCL_SUPPRESS_DEPRECATED_PUSH
SpmvParams<ValueT, OffsetT>& spmv_params;
_CCCL_SUPPRESS_DEPRECATED_POP

/// Wrapped pointer to the array of \p num_nonzeros values of the corresponding nonzero elements
/// of matrix <b>A</b>.
Expand Down Expand Up @@ -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<ValueT, OffsetT>& spmv_params)
: temp_storage(temp_storage.Alias())
, spmv_params(spmv_params)
Expand All @@ -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
Expand Down
31 changes: 18 additions & 13 deletions cub/cub/device/device_spmv.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
//! @{
Expand Down Expand Up @@ -177,22 +177,25 @@ 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");

_CCCL_SUPPRESS_DEPRECATED_PUSH
SpmvParams<ValueT, int> 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;
Expand All @@ -204,7 +207,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
}

#ifndef _CCCL_DOXYGEN_INVOKED // Do not document
Expand Down
33 changes: 30 additions & 3 deletions cub/cub/device/dispatch/dispatch_spmv_orig.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,7 @@

#include <nv/target>

_CCCL_SUPPRESS_DEPRECATED_PUSH
CUB_NAMESPACE_BEGIN

/******************************************************************************
Expand All @@ -83,7 +84,9 @@ CUB_NAMESPACE_BEGIN
* @param[in] spmv_params
* SpMV input parameter bundle
*/
_CCCL_SUPPRESS_DEPRECATED_PUSH
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 =
Expand All @@ -106,6 +109,7 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmv1ColKernel(SpmvParams<ValueT, Offset
spmv_params.d_vector_y[row_idx] = value;
}
}
_CCCL_SUPPRESS_DEPRECATED_POP

/**
* @brief Spmv search kernel. Identifies merge path starting coordinates for each tile.
Expand All @@ -132,8 +136,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
Expand Down Expand Up @@ -210,13 +215,15 @@ 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 <typename SpmvPolicyT,
typename ScanTileStateT,
typename ValueT,
typename OffsetT,
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,
Expand All @@ -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<SpmvPolicyT, ValueT, OffsetT, HAS_ALPHA, HAS_BETA>;
_CCCL_SUPPRESS_DEPRECATED_POP

// Shared memory for AgentSpmv
__shared__ typename AgentSpmvT::TempStorage temp_storage;
Expand All @@ -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
Expand All @@ -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 <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);
Expand All @@ -264,6 +276,7 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSpmvEmptyMatrixKernel(SpmvParams<ValueT,
spmv_params.d_vector_y[row] = result;
}
}
_CCCL_SUPPRESS_DEPRECATED_POP

/**
* @brief Multi-block reduce-by-key sweep kernel entry point
Expand Down Expand Up @@ -298,11 +311,13 @@ 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,
Expand All @@ -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
Expand All @@ -342,7 +358,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
Expand All @@ -355,7 +371,9 @@ struct DispatchSpmv
};

// SpmvParams bundle type
_CCCL_SUPPRESS_DEPRECATED_PUSH
using SpmvParamsT = SpmvParams<ValueT, OffsetT>;
_CCCL_SUPPRESS_DEPRECATED_POP

// 2D merge path coordinate type
using CoordinateT = typename CubVector<OffsetT, 2>::Type;
Expand All @@ -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,
Expand All @@ -383,13 +402,15 @@ 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>;
};

/// SM37
struct Policy370
{
_CCCL_SUPPRESS_DEPRECATED_PUSH
using SpmvPolicyT =
AgentSpmvPolicy<(sizeof(ValueT) > 4) ? 128 : 128,
(sizeof(ValueT) > 4) ? 9 : 14,
Expand All @@ -400,13 +421,15 @@ 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>;
};

/// SM50
struct Policy500
{
_CCCL_SUPPRESS_DEPRECATED_PUSH
using SpmvPolicyT =
AgentSpmvPolicy<(sizeof(ValueT) > 4) ? 64 : 128,
(sizeof(ValueT) > 4) ? 6 : 7,
Expand All @@ -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>;
Expand All @@ -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,
Expand All @@ -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>;
};
Expand Down Expand Up @@ -1005,4 +1031,5 @@ struct DispatchSpmv
#endif // _CCCL_DOXYGEN_INVOKED
};

_CCCL_SUPPRESS_DEPRECATED_POP
CUB_NAMESPACE_END
4 changes: 4 additions & 0 deletions cub/test/test_device_spmv.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,10 @@
#include <c2h/device_policy.h>
#include <c2h/vector.h>

// 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;

//==============================================================================
Expand Down

0 comments on commit d5ca93c

Please sign in to comment.