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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.6.0/hipcub/include/hipcub/backend/rocprim/grid/grid_queue.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/grid/grid_queue.hpp Source File

 /******************************************************************************
  * Copyright (c) 2011, Duane Merrill.  All rights reserved.
  * Copyright (c) 2011-2018, NVIDIA CORPORATION.  All rights reserved.
  * Modifications Copyright (c) 2021, 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_GRID_GRID_QUEUE_HPP_
 #define HIPCUB_ROCPRIM_GRID_GRID_QUEUE_HPP_
  
 #include <type_traits>
  
 #include "../../../config.hpp"
  
 BEGIN_HIPCUB_NAMESPACE
  
 template <typename OffsetT>
 class GridQueue
 {
 private:
  
     enum
     {
         FILL    = 0,
         DRAIN   = 1,
     };
  
     OffsetT *d_counters;
  
 public:
  
     __host__ __device__ __forceinline__
     static size_t AllocationSize()
     {
         return sizeof(OffsetT) * 2;
     }
  
  
     __host__ __device__ __forceinline__ GridQueue()
     :
         d_counters(NULL)
     {}
  
  
     __host__ __device__ __forceinline__ GridQueue(
         void *d_storage) 
     :
         d_counters((OffsetT*) d_storage)
     {}
  
  
     HIPCUB_DEVICE hipError_t FillAndResetDrain(
         OffsetT fill_size,
         hipStream_t stream = 0)
     {
         hipError_t result = hipErrorUnknown;
         (void)stream;
         d_counters[FILL] = fill_size;
         d_counters[DRAIN] = 0;
         result = hipSuccess;
         return result;
     }
  
     HIPCUB_HOST hipError_t FillAndResetDrain(
         OffsetT fill_size,
         hipStream_t stream = 0)
     {
         hipError_t result = hipErrorUnknown;
         OffsetT counters[2];
         counters[FILL] = fill_size;
         counters[DRAIN] = 0;
         result = CubDebug(hipMemcpyAsync(d_counters, counters, sizeof(OffsetT) * 2, hipMemcpyHostToDevice, stream));
         return result;
     }
  
     HIPCUB_DEVICE hipError_t ResetDrain(hipStream_t stream = 0)
     {
         hipError_t result = hipErrorUnknown;
         (void)stream;
         d_counters[DRAIN] = 0;
         result = hipSuccess;
         return result;
     }
  
     HIPCUB_HOST hipError_t ResetDrain(hipStream_t stream = 0)
     {
         hipError_t result = hipErrorUnknown;
         result = CubDebug(hipMemsetAsync(d_counters + DRAIN, 0, sizeof(OffsetT), stream));
         return result;
     }
  
  
     HIPCUB_DEVICE hipError_t ResetFill(hipStream_t stream = 0)
     {
         hipError_t result = hipErrorUnknown;
         (void)stream;
         d_counters[FILL] = 0;
         result = hipSuccess;
         return result;
     }
  
     HIPCUB_HOST hipError_t ResetFill(hipStream_t stream = 0)
     {
         hipError_t result = hipErrorUnknown;
         result = CubDebug(hipMemsetAsync(d_counters + FILL, 0, sizeof(OffsetT), stream));
         return result;
     }
  
  
     HIPCUB_DEVICE hipError_t FillSize(
         OffsetT &fill_size,
         hipStream_t stream = 0)
     {
         hipError_t result = hipErrorUnknown;
         (void)stream;
         fill_size = d_counters[FILL];
         result = hipSuccess;
         return result;
     }
  
     HIPCUB_HOST hipError_t FillSize(
         OffsetT &fill_size,
         hipStream_t stream = 0)
     {
         hipError_t result = hipErrorUnknown;
         result = CubDebug(hipMemcpyAsync(&fill_size, d_counters + FILL, sizeof(OffsetT), hipMemcpyDeviceToHost, stream));
         return result;
     }
  
  
     HIPCUB_DEVICE OffsetT Drain(OffsetT num_items)
     {
         return atomicAdd(d_counters + DRAIN, num_items);
     }
  
  
     HIPCUB_DEVICE OffsetT Fill(OffsetT num_items)
     {
         return atomicAdd(d_counters + FILL, num_items);
     }
 };
  
  
 #ifndef DOXYGEN_SHOULD_SKIP_THIS    // Do not document
  
  
 template <typename OffsetT>
 __global__ void FillAndResetDrainKernel(
     GridQueue<OffsetT>   grid_queue,
     OffsetT              num_items)
 {
     grid_queue.FillAndResetDrain(num_items);
 }
  
  
  
 #endif // DOXYGEN_SHOULD_SKIP_THIS
  
        // end group GridModule
  
 END_HIPCUB_NAMESPACE
  
 #endif // HIPCUB_ROCPRIM_GRID_GRID_QUEUE_HPP_