Extras/CUDA/radixsort.cu

   1 /*
   2  * Copyright 1993-2006 NVIDIA Corporation.  All rights reserved.
   3  *
   4  * NOTICE TO USER:
   5  *
   6  * This source code is subject to NVIDIA ownership rights under U.S. and
   7  * international Copyright laws.
   8  *
   9  * NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
  10  * CODE FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
  11  * IMPLIED WARRANTY OF ANY KIND.  NVIDIA DISCLAIMS ALL WARRANTIES WITH
  12  * REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
  13  * MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
  14  * IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,
  15  * OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS
  16  * OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE
  17  * OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE
  18  * OR PERFORMANCE OF THIS SOURCE CODE.
  19  *
  20  * U.S. Government End Users.  This source code is a "commercial item" as
  21  * that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting  of
  22  * "commercial computer software" and "commercial computer software
  23  * documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995)
  24  * and is provided to the U.S. Government only as a commercial end item.
  25  * Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
  26  * 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
  27  * source code with only those rights set forth herein.
  28  */
  29
  30 /* Radixsort project with key/value and arbitrary datset size support
  31  * which demonstrates the use of CUDA in a multi phase sorting
  32  * computation.
  33  * Host code.
  34  */
  35
  36 #include "radixsort.cuh"
  37 #include "radixsort_kernel.cu"
  38
  39 extern "C"
  40 {
  41
  42 ////////////////////////////////////////////////////////////////////////////////
  43 //! Perform a radix sort
  44 //! Sorting performed in place on passed arrays.
  45 //!
  46 //! @param pData0       input and output array - data will be sorted
  47 //! @param pData1       additional array to allow ping pong computation
  48 //! @param elements     number of elements to sort
  49 ////////////////////////////////////////////////////////////////////////////////
  50 void RadixSort(KeyValuePair *pData0, KeyValuePair *pData1, uint elements, uint bits)
  51 {
  52     // Round element count to total number of threads for efficiency
  53     uint elements_rounded_to_3072;
  54     int modval = elements % 3072;
  55     if( modval == 0 )
  56         elements_rounded_to_3072 = elements;
  57     else
  58         elements_rounded_to_3072 = elements + (3072 - (modval));
  59
  60     // Iterate over n bytes of y bit word, using each byte to sort the list in turn
  61     for (uint shift = 0; shift < bits; shift += RADIX)
  62     {
  63         // Perform one round of radix sorting
  64
  65         // Generate per radix group sums radix counts across a radix group
  66         RadixSum<<<NUM_BLOCKS, NUM_THREADS_PER_BLOCK, GRFSIZE>>>(pData0, elements, elements_rounded_to_3072, shift);
  67         // Prefix sum in radix groups, and then between groups throughout a block
  68         RadixPrefixSum<<<PREFIX_NUM_BLOCKS, PREFIX_NUM_THREADS_PER_BLOCK, PREFIX_GRFSIZE>>>();
  69         // Sum the block offsets and then shuffle data into bins
  70         RadixAddOffsetsAndShuffle<<<NUM_BLOCKS, NUM_THREADS_PER_BLOCK, SHUFFLE_GRFSIZE>>>(pData0, pData1, elements, elements_rounded_to_3072, shift);
  71
  72         // Exchange data pointers
  73         KeyValuePair* pTemp = pData0;
  74         pData0 = pData1;
  75         pData1 = pTemp;
  76    }
  77 }
  78
  79 }