Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Backport to 2.8: Deprecate cub::DeviceSpmv (#3320) #3374

Merged
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
15 changes: 12 additions & 3 deletions cub/cub/agent/agent_spmv_orig.cuh
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
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
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
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
Loading