block_scan_raking.cuh

Go to the documentation of this file.

/******************************************************************************
 * Copyright (c) 2011, Duane Merrill.  All rights reserved.
 * Copyright (c) 2011-2018, NVIDIA CORPORATION.  All rights reserved.
 * 
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions are met:
 *     * Redistributions of source code must retain the above copyright
 *       notice, this list of conditions and the following disclaimer.
 *     * Redistributions in binary form must reproduce the above copyright
 *       notice, this list of conditions and the following disclaimer in the
 *       documentation and/or other materials provided with the distribution.
 *     * Neither the name of the NVIDIA CORPORATION nor the
 *       names of its contributors may be used to endorse or promote products
 *       derived from this software without specific prior written permission.
 * 
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 *
 ******************************************************************************/


#pragma once

#include "../../util_ptx.cuh"
#include "../../util_arch.cuh"
#include "../../block/block_raking_layout.cuh"
#include "../../thread/thread_reduce.cuh"
#include "../../thread/thread_scan.cuh"
#include "../../warp/warp_scan.cuh"
#include "../../util_namespace.cuh"

CUB_NS_PREFIX

namespace cub {


template <
    typename    T,              
    int         BLOCK_DIM_X,    
    int         BLOCK_DIM_Y,    
    int         BLOCK_DIM_Z,    
    bool        MEMOIZE,        
    int         PTX_ARCH>       
struct BlockScanRaking
{
    //---------------------------------------------------------------------
    // Types and constants
    //---------------------------------------------------------------------

    enum
    {
        BLOCK_THREADS = BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z,
    };

    typedef BlockRakingLayout<T, BLOCK_THREADS, PTX_ARCH> BlockRakingLayout;

    enum
    {
        RAKING_THREADS = BlockRakingLayout::RAKING_THREADS,

        SEGMENT_LENGTH = BlockRakingLayout::SEGMENT_LENGTH,

        WARP_SYNCHRONOUS = (BLOCK_THREADS == RAKING_THREADS),
    };

    typedef WarpScan<T, RAKING_THREADS, PTX_ARCH> WarpScan;

    struct _TempStorage
    {
        typename WarpScan::TempStorage              warp_scan;          
        typename BlockRakingLayout::TempStorage     raking_grid;        
        T                                           block_aggregate;    
    };


    struct TempStorage : Uninitialized<_TempStorage> {};


    //---------------------------------------------------------------------
    // Per-thread fields
    //---------------------------------------------------------------------

    // Thread fields
    _TempStorage    &temp_storage;
    unsigned int    linear_tid;
    T               cached_segment[SEGMENT_LENGTH];


    //---------------------------------------------------------------------
    // Utility methods
    //---------------------------------------------------------------------

    template <int ITERATION, typename ScanOp>
    __device__ __forceinline__ T GuardedReduce(
        T*                  raking_ptr,         
        ScanOp              scan_op,            
        T                   raking_partial,     
        Int2Type<ITERATION> /*iteration*/)
    {
        if ((BlockRakingLayout::UNGUARDED) || (((linear_tid * SEGMENT_LENGTH) + ITERATION) < BLOCK_THREADS))
        {
            T addend = raking_ptr[ITERATION];
            raking_partial = scan_op(raking_partial, addend);
        }

        return GuardedReduce(raking_ptr, scan_op, raking_partial, Int2Type<ITERATION + 1>());
    }


    template <typename ScanOp>
    __device__ __forceinline__ T GuardedReduce(
        T*                          /*raking_ptr*/,    
        ScanOp                      /*scan_op*/,       
        T                           raking_partial,    
        Int2Type<SEGMENT_LENGTH>    /*iteration*/)
    {
        return raking_partial;
    }


    template <int ITERATION>
    __device__ __forceinline__ void CopySegment(
        T*                  out,            
        T*                  in,             
        Int2Type<ITERATION> /*iteration*/)
    {
        out[ITERATION] = in[ITERATION];
        CopySegment(out, in, Int2Type<ITERATION + 1>());
    }

 
    __device__ __forceinline__ void CopySegment(
        T*                  /*out*/,            
        T*                  /*in*/,             
        Int2Type<SEGMENT_LENGTH> /*iteration*/)
    {}


    template <typename ScanOp>
    __device__ __forceinline__ T Upsweep(
        ScanOp scan_op)
    {
        T *smem_raking_ptr = BlockRakingLayout::RakingPtr(temp_storage.raking_grid, linear_tid);

        // Read data into registers
        CopySegment(cached_segment, smem_raking_ptr, Int2Type<0>());

        T raking_partial = cached_segment[0];

        return GuardedReduce(cached_segment, scan_op, raking_partial, Int2Type<1>());
    }


    template <typename ScanOp>
    __device__ __forceinline__ void ExclusiveDownsweep(
        ScanOp          scan_op,
        T               raking_partial,
        bool            apply_prefix = true)
    {
        T *smem_raking_ptr = BlockRakingLayout::RakingPtr(temp_storage.raking_grid, linear_tid);

        // Read data back into registers
        if (!MEMOIZE)
        {
            CopySegment(cached_segment, smem_raking_ptr, Int2Type<0>());
        }

        internal::ThreadScanExclusive(cached_segment, cached_segment, scan_op, raking_partial, apply_prefix);

        // Write data back to smem
        CopySegment(smem_raking_ptr, cached_segment, Int2Type<0>());
    }


    template <typename ScanOp>
    __device__ __forceinline__ void InclusiveDownsweep(
        ScanOp          scan_op,
        T               raking_partial,
        bool            apply_prefix = true)
    {
        T *smem_raking_ptr = BlockRakingLayout::RakingPtr(temp_storage.raking_grid, linear_tid);

        // Read data back into registers
        if (!MEMOIZE)
        {
            CopySegment(cached_segment, smem_raking_ptr, Int2Type<0>());
        }

        internal::ThreadScanInclusive(cached_segment, cached_segment, scan_op, raking_partial, apply_prefix);

        // Write data back to smem
        CopySegment(smem_raking_ptr, cached_segment, Int2Type<0>());
    }


    //---------------------------------------------------------------------
    // Constructors
    //---------------------------------------------------------------------

    __device__ __forceinline__ BlockScanRaking(
        TempStorage &temp_storage)
    :
        temp_storage(temp_storage.Alias()),
        linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z))
    {}


    //---------------------------------------------------------------------
    // Exclusive scans
    //---------------------------------------------------------------------

    template <typename ScanOp>
    __device__ __forceinline__ void ExclusiveScan(
        T               input,                          
        T               &exclusive_output,                        
        ScanOp          scan_op)                        
    {
        if (WARP_SYNCHRONOUS)
        {
            // Short-circuit directly to warp-synchronous scan
            WarpScan(temp_storage.warp_scan).ExclusiveScan(input, exclusive_output, scan_op);
        }
        else
        {
            // Place thread partial into shared memory raking grid
            T *placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
            *placement_ptr = input;

            CTA_SYNC();

            // Reduce parallelism down to just raking threads
            if (linear_tid < RAKING_THREADS)
            {
                // Raking upsweep reduction across shared partials
                T upsweep_partial = Upsweep(scan_op);

                // Warp-synchronous scan
                T exclusive_partial;
                WarpScan(temp_storage.warp_scan).ExclusiveScan(upsweep_partial, exclusive_partial, scan_op);

                // Exclusive raking downsweep scan
                ExclusiveDownsweep(scan_op, exclusive_partial, (linear_tid != 0));
            }

            CTA_SYNC();

            // Grab thread prefix from shared memory
            exclusive_output = *placement_ptr;
        }
    }

    template <typename ScanOp>
    __device__ __forceinline__ void ExclusiveScan(
        T               input,              
        T               &output,            
        const T         &initial_value,     
        ScanOp          scan_op)            
    {
        if (WARP_SYNCHRONOUS)
        {
            // Short-circuit directly to warp-synchronous scan
            WarpScan(temp_storage.warp_scan).ExclusiveScan(input, output, initial_value, scan_op);
        }
        else
        {
            // Place thread partial into shared memory raking grid
            T *placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
            *placement_ptr = input;

            CTA_SYNC();

            // Reduce parallelism down to just raking threads
            if (linear_tid < RAKING_THREADS)
            {
                // Raking upsweep reduction across shared partials
                T upsweep_partial = Upsweep(scan_op);

                // Exclusive Warp-synchronous scan
                T exclusive_partial;
                WarpScan(temp_storage.warp_scan).ExclusiveScan(upsweep_partial, exclusive_partial, initial_value, scan_op);

                // Exclusive raking downsweep scan
                ExclusiveDownsweep(scan_op, exclusive_partial);
            }

            CTA_SYNC();

            // Grab exclusive partial from shared memory
            output = *placement_ptr;
        }
    }


    template <typename ScanOp>
    __device__ __forceinline__ void ExclusiveScan(
        T               input,                          
        T               &output,                        
        ScanOp          scan_op,                        
        T               &block_aggregate)               
    {
        if (WARP_SYNCHRONOUS)
        {
            // Short-circuit directly to warp-synchronous scan
            WarpScan(temp_storage.warp_scan).ExclusiveScan(input, output, scan_op, block_aggregate);
        }
        else
        {
            // Place thread partial into shared memory raking grid
            T *placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
            *placement_ptr = input;

            CTA_SYNC();

            // Reduce parallelism down to just raking threads
            if (linear_tid < RAKING_THREADS)
            {
                // Raking upsweep reduction across shared partials
                T upsweep_partial= Upsweep(scan_op);

                // Warp-synchronous scan
                T inclusive_partial;
                T exclusive_partial;
                WarpScan(temp_storage.warp_scan).Scan(upsweep_partial, inclusive_partial, exclusive_partial, scan_op);

                // Exclusive raking downsweep scan
                ExclusiveDownsweep(scan_op, exclusive_partial, (linear_tid != 0));

                // Broadcast aggregate to all threads
                if (linear_tid == RAKING_THREADS - 1)
                    temp_storage.block_aggregate = inclusive_partial;
            }

            CTA_SYNC();

            // Grab thread prefix from shared memory
            output = *placement_ptr;

            // Retrieve block aggregate
            block_aggregate = temp_storage.block_aggregate;
        }
    }


    template <typename ScanOp>
    __device__ __forceinline__ void ExclusiveScan(
        T               input,              
        T               &output,            
        const T         &initial_value,     
        ScanOp          scan_op,            
        T               &block_aggregate)   
    {
        if (WARP_SYNCHRONOUS)
        {
            // Short-circuit directly to warp-synchronous scan
            WarpScan(temp_storage.warp_scan).ExclusiveScan(input, output, initial_value, scan_op, block_aggregate);
        }
        else
        {
            // Place thread partial into shared memory raking grid
            T *placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
            *placement_ptr = input;

            CTA_SYNC();

            // Reduce parallelism down to just raking threads
            if (linear_tid < RAKING_THREADS)
            {
                // Raking upsweep reduction across shared partials
                T upsweep_partial = Upsweep(scan_op);

                // Warp-synchronous scan
                T exclusive_partial;
                WarpScan(temp_storage.warp_scan).ExclusiveScan(upsweep_partial, exclusive_partial, initial_value, scan_op, block_aggregate);

                // Exclusive raking downsweep scan
                ExclusiveDownsweep(scan_op, exclusive_partial);

                // Broadcast aggregate to other threads
                if (linear_tid == 0)
                    temp_storage.block_aggregate = block_aggregate;
            }

            CTA_SYNC();

            // Grab exclusive partial from shared memory
            output = *placement_ptr;

            // Retrieve block aggregate
            block_aggregate = temp_storage.block_aggregate;
        }
    }


    template <
        typename ScanOp,
        typename BlockPrefixCallbackOp>
    __device__ __forceinline__ void ExclusiveScan(
        T                       input,                          
        T                       &output,                        
        ScanOp                  scan_op,                        
        BlockPrefixCallbackOp   &block_prefix_callback_op)      
    {
        if (WARP_SYNCHRONOUS)
        {
            // Short-circuit directly to warp-synchronous scan
            T block_aggregate;
            WarpScan warp_scan(temp_storage.warp_scan);
            warp_scan.ExclusiveScan(input, output, scan_op, block_aggregate);

            // Obtain warp-wide prefix in lane0, then broadcast to other lanes
            T block_prefix = block_prefix_callback_op(block_aggregate);
            block_prefix = warp_scan.Broadcast(block_prefix, 0);

            output = scan_op(block_prefix, output);
            if (linear_tid == 0)
                output = block_prefix;
        }
        else
        {
            // Place thread partial into shared memory raking grid
            T *placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
            *placement_ptr = input;

            CTA_SYNC();

            // Reduce parallelism down to just raking threads
            if (linear_tid < RAKING_THREADS)
            {
                WarpScan warp_scan(temp_storage.warp_scan);

                // Raking upsweep reduction across shared partials
                T upsweep_partial = Upsweep(scan_op);

                // Warp-synchronous scan
                T exclusive_partial, block_aggregate;
                warp_scan.ExclusiveScan(upsweep_partial, exclusive_partial, scan_op, block_aggregate);

                // Obtain block-wide prefix in lane0, then broadcast to other lanes
                T block_prefix = block_prefix_callback_op(block_aggregate);
                block_prefix = warp_scan.Broadcast(block_prefix, 0);

                // Update prefix with warpscan exclusive partial
                T downsweep_prefix = scan_op(block_prefix, exclusive_partial);
                if (linear_tid == 0)
                    downsweep_prefix = block_prefix;

                // Exclusive raking downsweep scan
                ExclusiveDownsweep(scan_op, downsweep_prefix);
            }

            CTA_SYNC();

            // Grab thread prefix from shared memory
            output = *placement_ptr;
        }
    }


    //---------------------------------------------------------------------
    // Inclusive scans
    //---------------------------------------------------------------------

    template <typename ScanOp>
    __device__ __forceinline__ void InclusiveScan(
        T               input,                          
        T               &output,                        
        ScanOp          scan_op)                        
    {
        if (WARP_SYNCHRONOUS)
        {
            // Short-circuit directly to warp-synchronous scan
            WarpScan(temp_storage.warp_scan).InclusiveScan(input, output, scan_op);
        }
        else
        {
            // Place thread partial into shared memory raking grid
            T *placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
            *placement_ptr = input;

            CTA_SYNC();

            // Reduce parallelism down to just raking threads
            if (linear_tid < RAKING_THREADS)
            {
                // Raking upsweep reduction across shared partials
                T upsweep_partial = Upsweep(scan_op);

                // Exclusive Warp-synchronous scan
                T exclusive_partial;
                WarpScan(temp_storage.warp_scan).ExclusiveScan(upsweep_partial, exclusive_partial, scan_op);

                // Inclusive raking downsweep scan
                InclusiveDownsweep(scan_op, exclusive_partial, (linear_tid != 0));
            }

            CTA_SYNC();

            // Grab thread prefix from shared memory
            output = *placement_ptr;
        }
    }


    template <typename ScanOp>
    __device__ __forceinline__ void InclusiveScan(
        T               input,                          
        T               &output,                        
        ScanOp          scan_op,                        
        T               &block_aggregate)               
    {
        if (WARP_SYNCHRONOUS)
        {
            // Short-circuit directly to warp-synchronous scan
            WarpScan(temp_storage.warp_scan).InclusiveScan(input, output, scan_op, block_aggregate);
        }
        else
        {
            // Place thread partial into shared memory raking grid
            T *placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
            *placement_ptr = input;

            CTA_SYNC();

            // Reduce parallelism down to just raking threads
            if (linear_tid < RAKING_THREADS)
            {
                // Raking upsweep reduction across shared partials
                T upsweep_partial = Upsweep(scan_op);

                // Warp-synchronous scan
                T inclusive_partial;
                T exclusive_partial;
                WarpScan(temp_storage.warp_scan).Scan(upsweep_partial, inclusive_partial, exclusive_partial, scan_op);

                // Inclusive raking downsweep scan
                InclusiveDownsweep(scan_op, exclusive_partial, (linear_tid != 0));

                // Broadcast aggregate to all threads
                if (linear_tid == RAKING_THREADS - 1)
                    temp_storage.block_aggregate = inclusive_partial;
            }

            CTA_SYNC();

            // Grab thread prefix from shared memory
            output = *placement_ptr;

            // Retrieve block aggregate
            block_aggregate = temp_storage.block_aggregate;
        }
    }


    template <
        typename ScanOp,
        typename BlockPrefixCallbackOp>
    __device__ __forceinline__ void InclusiveScan(
        T                       input,                          
        T                       &output,                        
        ScanOp                  scan_op,                        
        BlockPrefixCallbackOp   &block_prefix_callback_op)      
    {
        if (WARP_SYNCHRONOUS)
        {
            // Short-circuit directly to warp-synchronous scan
            T block_aggregate;
            WarpScan warp_scan(temp_storage.warp_scan);
            warp_scan.InclusiveScan(input, output, scan_op, block_aggregate);

            // Obtain warp-wide prefix in lane0, then broadcast to other lanes
            T block_prefix = block_prefix_callback_op(block_aggregate);
            block_prefix = warp_scan.Broadcast(block_prefix, 0);

            // Update prefix with exclusive warpscan partial
            output = scan_op(block_prefix, output);
        }
        else
        {
            // Place thread partial into shared memory raking grid
            T *placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
            *placement_ptr = input;

            CTA_SYNC();

            // Reduce parallelism down to just raking threads
            if (linear_tid < RAKING_THREADS)
            {
                WarpScan warp_scan(temp_storage.warp_scan);

                // Raking upsweep reduction across shared partials
                T upsweep_partial = Upsweep(scan_op);

                // Warp-synchronous scan
                T exclusive_partial, block_aggregate;
                warp_scan.ExclusiveScan(upsweep_partial, exclusive_partial, scan_op, block_aggregate);

                // Obtain block-wide prefix in lane0, then broadcast to other lanes
                T block_prefix = block_prefix_callback_op(block_aggregate);
                block_prefix = warp_scan.Broadcast(block_prefix, 0);

                // Update prefix with warpscan exclusive partial
                T downsweep_prefix = scan_op(block_prefix, exclusive_partial);
                if (linear_tid == 0)
                    downsweep_prefix = block_prefix;

                // Inclusive raking downsweep scan
                InclusiveDownsweep(scan_op, downsweep_prefix);
            }

            CTA_SYNC();

            // Grab thread prefix from shared memory
            output = *placement_ptr;
        }
    }

};


}               // CUB namespace
CUB_NS_POSTFIX  // Optional outer namespace(s)