forked from lijiext/lammps
114 lines
4.2 KiB
Plaintext
114 lines
4.2 KiB
Plaintext
// -------------------------------------------------------------
|
|
// cuDPP -- CUDA Data Parallel Primitives library
|
|
// -------------------------------------------------------------
|
|
// $Revision: 5633 $
|
|
// $Date: 2009-07-01 15:02:51 +1000 (Wed, 01 Jul 2009) $
|
|
// -------------------------------------------------------------
|
|
// This source code is distributed under the terms of license.txt
|
|
// in the root directory of this source distribution.
|
|
// -------------------------------------------------------------
|
|
|
|
/**
|
|
* @file
|
|
* scan_kernel.cu
|
|
*
|
|
* @brief CUDPP kernel-level scan routines
|
|
*/
|
|
|
|
/** \defgroup cudpp_kernel CUDPP Kernel-Level API
|
|
* The CUDPP Kernel-Level API contains functions that run on the GPU
|
|
* device across a grid of Cooperative Thread Array (CTA, aka Thread
|
|
* Block). These kernels are declared \c __global__ so that they
|
|
* must be invoked from host (CPU) code. They generally invoke GPU
|
|
* \c __device__ routines in the CUDPP \link cudpp_cta CTA-Level API\endlink.
|
|
* Kernel-Level API functions are used by CUDPP
|
|
* \link cudpp_app Application-Level\endlink functions to implement their
|
|
* functionality.
|
|
* @{
|
|
*/
|
|
|
|
/** @name Scan Functions
|
|
* @{
|
|
*/
|
|
|
|
#include <cudpp_globals.h>
|
|
#include "cta/scan_cta.cu"
|
|
#include "sharedmem.h"
|
|
|
|
/**
|
|
* @brief Main scan kernel
|
|
*
|
|
* This __global__ device function performs one level of a multiblock scan on
|
|
* an arbitrary-dimensioned array in \a d_in, returning the result in \a d_out
|
|
* (which may point to the same array). The same function may be used for
|
|
* single or multi-row scans. To perform a multirow scan, pass the width of
|
|
* each row of the input row (in elements) in \a dataRowPitch, and the width of
|
|
* the rows of \a d_blockSums (in elements) in \a blockSumRowPitch, and invoke
|
|
* with a thread block grid with height greater than 1.
|
|
*
|
|
* This function peforms one level of a recursive, multiblock scan. At the
|
|
* app level, this function is called by cudppScan and cudppMultiScan and used
|
|
* in combination with vectorAddUniform4() to produce a complete scan.
|
|
*
|
|
* Template parameter \a T is the datatype of the array to be scanned.
|
|
* Template parameter \a traits is the ScanTraits struct containing
|
|
* compile-time options for the scan, such as whether it is forward or
|
|
* backward, exclusive or inclusive, multi- or single-row, etc.
|
|
*
|
|
* @param[out] d_out The output (scanned) array
|
|
* @param[in] d_in The input array to be scanned
|
|
* @param[out] d_blockSums The array of per-block sums
|
|
* @param[in] numElements The number of elements to scan
|
|
* @param[in] dataRowPitch The width of each row of \a d_in in elements
|
|
* (for multi-row scans)
|
|
* @param[in] blockSumRowPitch The with of each row of \a d_blockSums in elements
|
|
* (for multi-row scans)
|
|
*/
|
|
template<class T, class traits>
|
|
__global__ void scan4(T *d_out,
|
|
const T *d_in,
|
|
T *d_blockSums,
|
|
int numElements,
|
|
unsigned int dataRowPitch,
|
|
unsigned int blockSumRowPitch)
|
|
{
|
|
SharedMemory<T> smem;
|
|
T* temp = smem.getPointer();
|
|
|
|
int devOffset, ai, bi, aiDev, biDev;
|
|
T threadScan0[4], threadScan1[4];
|
|
|
|
unsigned int blockN = numElements;
|
|
unsigned int blockSumIndex = blockIdx.x;
|
|
|
|
if (traits::isMultiRow())
|
|
{
|
|
//int width = __mul24(gridDim.x, blockDim.x) << 1;
|
|
int yIndex = __umul24(blockDim.y, blockIdx.y) + threadIdx.y;
|
|
devOffset = __umul24(dataRowPitch, yIndex);
|
|
blockN += (devOffset << 2);
|
|
devOffset += __umul24(blockIdx.x, blockDim.x << 1);
|
|
blockSumIndex += __umul24(blockSumRowPitch << 2, yIndex) ;
|
|
}
|
|
else
|
|
{
|
|
devOffset = __umul24(blockIdx.x, (blockDim.x << 1));
|
|
}
|
|
|
|
// load data into shared memory
|
|
loadSharedChunkFromMem4<T, traits>
|
|
(temp, threadScan0, threadScan1, d_in,
|
|
blockN, devOffset, ai, bi, aiDev, biDev);
|
|
|
|
scanCTA<T, traits>(temp, d_blockSums, blockSumIndex);
|
|
|
|
// write results to device memory
|
|
storeSharedChunkToMem4<T, traits>
|
|
(d_out, threadScan0, threadScan1, temp,
|
|
blockN, devOffset, ai, bi, aiDev, biDev);
|
|
|
|
}
|
|
|
|
/** @} */ // end scan functions
|
|
/** @} */ // end cudpp_kernel
|