Skip to content

Commit

Permalink
- Fix initialization of the std::array in the CUB attribute cache.
Browse files Browse the repository at this point in the history
- Change C++11 type aliases in the refactored scan dispatch layer to be
  typedefs instead for C++03 compatibility.

Reviewed by: Allison Vacanti <[email protected]>
Reviewed by: Michał 'Griwes' Dominiak <[email protected]>
  • Loading branch information
brycelelbach committed Mar 30, 2020
1 parent 7513f31 commit 5102543
Show file tree
Hide file tree
Showing 2 changed files with 9 additions and 8 deletions.
15 changes: 8 additions & 7 deletions cub/device/dispatch/dispatch_scan.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -105,7 +105,7 @@ __global__ void DeviceScanKernel(
OutputIteratorT d_out, ///< Output data
ScanTileStateT tile_state, ///< Tile status interface
int start_tile, ///< The starting tile for the current grid
ScanOpT scan_op, ///< Binary scan functor
ScanOpT scan_op, ///< Binary scan functor
InitValueT init_value, ///< Initial value to seed the exclusive scan
OffsetT num_items) ///< Total number of scan items for the entire problem
{
Expand Down Expand Up @@ -308,7 +308,7 @@ struct DispatchScan:
debug_synchronous(debug_synchronous),
ptx_version(ptx_version)
{}

template <typename ActivePolicyT, typename InitKernel, typename ScanKernel>
CUB_RUNTIME_FUNCTION __host__ __forceinline__
cudaError_t Invoke(InitKernel init_kernel, ScanKernel scan_kernel)
Expand All @@ -323,8 +323,9 @@ struct DispatchScan:

#else

using Policy = typename ActivePolicyT::ScanPolicyT;
using ScanTileStateT = typename cub::ScanTileState<OutputT>;
typedef typename ActivePolicyT::ScanPolicyT Policy;
typedef typename cub::ScanTileState<OutputT> ScanTileStateT;

cudaError error = cudaSuccess;
do
{
Expand Down Expand Up @@ -424,8 +425,8 @@ struct DispatchScan:
CUB_RUNTIME_FUNCTION __host__ __forceinline__
cudaError_t Invoke()
{
using Policy = typename ActivePolicyT::ScanPolicyT;
using ScanTileStateT = typename cub::ScanTileState<OutputT>;
typedef typename ActivePolicyT::ScanPolicyT Policy;
typedef typename cub::ScanTileState<OutputT> ScanTileStateT;
// Ensure kernels are instantiated.
return Invoke<ActivePolicyT>(
DeviceScanInitKernel<ScanTileStateT>,
Expand All @@ -452,7 +453,7 @@ struct DispatchScan:
typedef typename DispatchScan::MaxPolicy MaxPolicyT;

cudaError_t error;
do
do
{
// Get PTX version
int ptx_version = 0;
Expand Down
2 changes: 1 addition & 1 deletion cub/util_device.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -263,7 +263,7 @@ public:
/**
* \brief Construct the cache.
*/
__host__ __forceinline__ PerDeviceAttributeCache() : entries_{}
__host__ __forceinline__ PerDeviceAttributeCache() : entries_()
{
assert(DeviceCount() <= CUB_MAX_DEVICES);
}
Expand Down

0 comments on commit 5102543

Please sign in to comment.