From 5102543d61b5ca67e1cfdb34d416feef7cf8eb57 Mon Sep 17 00:00:00 2001 From: Bryce Adelstein Lelbach aka wash Date: Mon, 30 Mar 2020 13:23:44 -0700 Subject: [PATCH] - Fix initialization of the `std::array` in the CUB attribute cache. - Change C++11 type aliases in the refactored scan dispatch layer to be typedefs instead for C++03 compatibility. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reviewed by: Allison Vacanti Reviewed by: Michał 'Griwes' Dominiak --- cub/device/dispatch/dispatch_scan.cuh | 15 ++++++++------- cub/util_device.cuh | 2 +- 2 files changed, 9 insertions(+), 8 deletions(-) diff --git a/cub/device/dispatch/dispatch_scan.cuh b/cub/device/dispatch/dispatch_scan.cuh index 23f47cd60b..667c631cd0 100644 --- a/cub/device/dispatch/dispatch_scan.cuh +++ b/cub/device/dispatch/dispatch_scan.cuh @@ -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 { @@ -308,7 +308,7 @@ struct DispatchScan: debug_synchronous(debug_synchronous), ptx_version(ptx_version) {} - + template CUB_RUNTIME_FUNCTION __host__ __forceinline__ cudaError_t Invoke(InitKernel init_kernel, ScanKernel scan_kernel) @@ -323,8 +323,9 @@ struct DispatchScan: #else - using Policy = typename ActivePolicyT::ScanPolicyT; - using ScanTileStateT = typename cub::ScanTileState; + typedef typename ActivePolicyT::ScanPolicyT Policy; + typedef typename cub::ScanTileState ScanTileStateT; + cudaError error = cudaSuccess; do { @@ -424,8 +425,8 @@ struct DispatchScan: CUB_RUNTIME_FUNCTION __host__ __forceinline__ cudaError_t Invoke() { - using Policy = typename ActivePolicyT::ScanPolicyT; - using ScanTileStateT = typename cub::ScanTileState; + typedef typename ActivePolicyT::ScanPolicyT Policy; + typedef typename cub::ScanTileState ScanTileStateT; // Ensure kernels are instantiated. return Invoke( DeviceScanInitKernel, @@ -452,7 +453,7 @@ struct DispatchScan: typedef typename DispatchScan::MaxPolicy MaxPolicyT; cudaError_t error; - do + do { // Get PTX version int ptx_version = 0; diff --git a/cub/util_device.cuh b/cub/util_device.cuh index 77128a3f0f..bfef8b7a76 100644 --- a/cub/util_device.cuh +++ b/cub/util_device.cuh @@ -263,7 +263,7 @@ public: /** * \brief Construct the cache. */ - __host__ __forceinline__ PerDeviceAttributeCache() : entries_{} + __host__ __forceinline__ PerDeviceAttributeCache() : entries_() { assert(DeviceCount() <= CUB_MAX_DEVICES); }