<commit_before>// Copyright (c) 2009-2017 The Regents of the University of Michigan
// This file is part of the HOOMD-blue project, released under the BSD 3-Clause License.

// Maintainer: mphoward

/*!
 * \file mpcd/CellThermoComputeGPU.cu
 * \brief Explicitly instantiates reduction operators and declares kernel drivers
 *        for mpcd::CellThermoComputeGPU.
 */

#include "CellThermoComputeGPU.cuh"
#include "CellThermoTypes.h"

#include "CellCommunicator.cuh"
#include "ReductionOperators.h"

#include "hoomd/extern/cub/cub/cub.cuh"

namespace mpcd
{
namespace gpu
{

//! CTA reduce, returns result in first thread
template<int LOGICAL_WARP_SIZE, typename T>
__device__ static T warp_reduce(T val)
    {
    static_assert(LOGICAL_WARP_SIZE <= CUB_PTX_WARP_THREADS, "Logical warp size cannot exceed hardware warp size");
    static_assert(LOGICAL_WARP_SIZE && !(LOGICAL_WARP_SIZE & (LOGICAL_WARP_SIZE-1)), "Logical warp size must be a power of 2");

    #pragma unroll
    for (int dest_count = LOGICAL_WARP_SIZE/2; dest_count >= 1; dest_count /= 2)
        {
        val += cub::ShuffleDown(val, dest_count);
        }
    return val;
    }

namespace kernel
{
//! Begins the cell thermo compute by summing cell quantities
/*!
 * \param d_cell_vel Velocity and mass per cell (output)
 * \param d_cell_energy Energy, temperature, number of particles per cell (output)
 * \param d_cell_np Number of particles per cell
 * \param d_cell_list MPCD cell list
 * \param cli Indexer into the cell list
 * \param d_vel MPCD particle velocities
 * \param N_mpcd Number of MPCD particles
 * \param mpcd_mass Mass of MPCD particle<commit_msg>Add documentation for warp reduction<commit_after>// Copyright (c) 2009-2017 The Regents of the University of Michigan
// This file is part of the HOOMD-blue project, released under the BSD 3-Clause License.

// Maintainer: mphoward

/*!
 * \file mpcd/CellThermoComputeGPU.cu
 * \brief Explicitly instantiates reduction operators and declares kernel drivers
 *        for mpcd::CellThermoComputeGPU.
 */

#include "CellThermoComputeGPU.cuh"
#include "CellThermoTypes.h"

#include "CellCommunicator.cuh"
#include "ReductionOperators.h"

#include "hoomd/extern/cub/cub/cub.cuh"

namespace mpcd
{
namespace gpu
{

//! Shuffle-based warp reduction
/*!
 * \param val Value to be reduced
 *
 * \tparam LOGICAL_WARP_SIZE Number of threads in a "logical" warp to reduce, must be a power-of-two
 *                           and less than the hardware warp size.
 * \tparam T Type of value to be reduced (inferred).
 *
 * \returns Reduced value.
 *
 * The value \a val is reduced into the 0-th lane of the "logical" warp using
 * shuffle-based intrinsics. This allows for the summation of quantities when
 * using multiple threads per object within a kernel.
 */
template<int LOGICAL_WARP_SIZE, typename T>
__device__ static T warp_reduce(T val)
    {
    static_assert(LOGICAL_WARP_SIZE <= CUB_PTX_WARP_THREADS, "Logical warp size cannot exceed hardware warp size");
    static_assert(LOGICAL_WARP_SIZE && !(LOGICAL_WARP_SIZE & (LOGICAL_WARP_SIZE-1)), "Logical warp size must be a power of 2");

    #pragma unroll
    for (int dest_count = LOGICAL_WARP_SIZE/2; dest_count >= 1; dest_count /= 2)
        {
        val += cub::ShuffleDown(val, dest_count);
        }
    return val;
    }

namespace kernel
{
//! Begins the cell thermo compute by summing cell quantities
/*!
 * \param d_cell_vel Velocity and mass per cell (output)
 * \param d_cell_energy Energy, temperature, number of particles per cell (output)
 * \param d_cell_np Number of particles per cell
 * \param d_cell_list MPCD cell list
 * \param cli Indexer into the cell list
 * \param d_vel MPCD particle velocities
 * \param N_mpcd Number of MPCD particles
 * \param mpcd_mass Mass of MPCD particle
{'fraction_non_alphanumeric': 0.08513442277279916, 'fraction_numerical': 0.007643647865050079, 'mean_word_length': 3.7916666666666665, 'pattern_counts': {'":': 0, '<': 7, '<?xml version=': 0, '>': 7, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 2, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '7371747', 'n_tokens_mistral': 1256, 'n_tokens_neox': 1179, 'n_words': 418}