/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.6.0/hipcub/include/hipcub/backend/rocprim/block/block_shuffle.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.6.0/hipcub/include/hipcub/backend/rocprim/block/block_shuffle.hpp Source File#

hipCUB: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.6.0/hipcub/include/hipcub/backend/rocprim/block/block_shuffle.hpp Source File

 /******************************************************************************
  * Copyright (c) 2010-2011, Duane Merrill.  All rights reserved.
  * Copyright (c) 2011-2018, NVIDIA CORPORATION.  All rights reserved.
  * Modifications Copyright (c) 2017-2020, Advanced Micro Devices, Inc.  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.
  *
  ******************************************************************************/
  
 #ifndef HIPCUB_ROCPRIM_BLOCK_BLOCK_SHUFFLE_HPP_
 #define HIPCUB_ROCPRIM_BLOCK_BLOCK_SHUFFLE_HPP_
  
 #include <type_traits>
  
 #include "../../../config.hpp"
  
 #include "../thread/thread_operators.hpp"
  
 #include <rocprim/block/block_shuffle.hpp>
  
 BEGIN_HIPCUB_NAMESPACE
  
  
  
 template <
     typename            T,
     int                 BLOCK_DIM_X,
     int                 BLOCK_DIM_Y         = 1,
     int                 BLOCK_DIM_Z         = 1,
     int                 ARCH            = HIPCUB_ARCH>
 class BlockShuffle : public ::rocprim::block_shuffle<
                     T,
                     BLOCK_DIM_X,
                     BLOCK_DIM_Y,
                     BLOCK_DIM_Z>
 {
   static_assert(
       BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z > 0,
       "BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z must be greater than 0"
   );
  
   using base_type =
       typename ::rocprim::block_shuffle<
           T,
           BLOCK_DIM_X,
           BLOCK_DIM_Y,
           BLOCK_DIM_Z
       >;
  
   // Reference to temporary storage (usually shared memory)
   typename base_type::storage_type& temp_storage_;
  
 public:
   using TempStorage = typename base_type::storage_type;
  
   HIPCUB_DEVICE inline
   BlockShuffle()  :      temp_storage_(private_storage())
   {}
  
  
   HIPCUB_DEVICE inline
   BlockShuffle(TempStorage &temp_storage)             
   :      temp_storage_(temp_storage)
   {}
  
   HIPCUB_DEVICE inline void Offset(
       T   input,                  
       T&  output,                 
       int distance = 1)           
   {
     base_type::offset(input,output,distance);
   }
  
   HIPCUB_DEVICE inline void Rotate(
       T   input,                  
       T&  output,                 
       unsigned int distance = 1)  
   {
     base_type::rotate(input,output,distance);
   }
   template <int ITEMS_PER_THREAD>
   HIPCUB_DEVICE inline void Up(
     T (&input)[ITEMS_PER_THREAD],   
     T (&prev)[ITEMS_PER_THREAD])    
   {
     base_type::up(input,prev);
   }
  
  
   template <int ITEMS_PER_THREAD>
   HIPCUB_DEVICE inline void Up(
       T (&input)[ITEMS_PER_THREAD],   
       T (&prev)[ITEMS_PER_THREAD],    
       T &block_suffix)                
   {
     base_type::up(input,prev,block_suffix);
   }
  
   template <int ITEMS_PER_THREAD>
   HIPCUB_DEVICE inline void Down(
       T (&input)[ITEMS_PER_THREAD],   
       T (&next)[ITEMS_PER_THREAD])    
   {
     base_type::down(input,next);
   }
  
   template <int ITEMS_PER_THREAD>
   HIPCUB_DEVICE inline void Down(
       T (&input)[ITEMS_PER_THREAD],   
       T (&next)[ITEMS_PER_THREAD],    
       T &block_prefix)                
   {
     base_type::down(input,next,block_prefix);
   }
  
 private:
     HIPCUB_DEVICE inline
     TempStorage& private_storage()
     {
         HIPCUB_SHARED_MEMORY TempStorage private_storage;
         return private_storage;
     }
 };
  
 END_HIPCUB_NAMESPACE
  
 #endif // HIPCUB_ROCPRIM_BLOCK_BLOCK_SHUFFLE_HPP_