← Back to branch summary

~ubuntu-branches/ubuntu/precise/openwalnut/precise

« back to all changes in this revision

Viewing changes to src/modules/detTractClustering/WCudaReduce.cu

  • Committer: Bazaar Package Importer
  • Author(s): Sebastian Eichelbaum
  • Date: 2011-06-21 10:26:54 UTC
  • Revision ID: james.westby@ubuntu.com-20110621102654-rq0zf436q949biih
Tags: upstream-1.2.5
Import upstream version 1.2.5

Show diffs side-by-side

added added

removed removed

Lines of Context:
src/modules/detTractClustering/WCudaReduce.cu
 
1
//---------------------------------------------------------------------------
 
2
//
 
3
// Project: OpenWalnut ( http://www.openwalnut.org )
 
4
//
 
5
// Copyright 2009 OpenWalnut Community, BSV@Uni-Leipzig and CNCF@MPI-CBS, Copyright 2010 RRZK University of Cologne
 
6
// For more information see http://www.openwalnut.org/copying
 
7
//
 
8
// This file is part of OpenWalnut.
 
9
//
 
10
// OpenWalnut is free software: you can redistribute it and/or modify
 
11
// it under the terms of the GNU Lesser General Public License as published by
 
12
// the Free Software Foundation, either version 3 of the License, or
 
13
// (at your option) any later version.
 
14
//
 
15
// OpenWalnut is distributed in the hope that it will be useful,
 
16
// but WITHOUT ANY WARRANTY; without even the implied warranty of
 
17
// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
 
18
// GNU Lesser General Public License for more details.
 
19
//
 
20
// You should have received a copy of the GNU Lesser General Public License
 
21
// along with OpenWalnut. If not, see <http://www.gnu.org/licenses/>.
 
22
//
 
23
//---------------------------------------------------------------------------
 
24
 
 
25
 
 
26
// NOTE: This is based on code from projects/reduction/reduction_kernel.cu from CUDA SDK 2.0
 
27
 
 
28
#ifndef WCUDAREDUCE_CU
 
29
#define WCUDAREDUCE_CU
 
30
 
 
31
/**
 
32
 * Performs parallel reduction of elements in \e shared memory, the reduction operation is defined by \c Op::eval.
 
33
 */
 
34
template< class Op >
 
35
class Reduce
 
36
{
 
37
public:
 
38
    /**
 
39
     * Result of the reduction is stored in s[0]
 
40
     */
 
41
    static inline __device__ void apply( volatile float *s, unsigned int n )
 
42
    {
 
43
        const unsigned int i = threadIdx.x;
 
44
 
 
45
        const int nthreads = min( 512, 1<<(31-__clz(min(blockDim.x, n))) );
 
46
        if( n > nthreads )
 
47
        {
 
48
            if( i < nthreads )
 
49
            {
 
50
                for( unsigned int j = nthreads+i; j<n; j += nthreads )
 
51
                {
 
52
                    s[i] = Op::eval( s[i], s[j] );
 
53
                }
 
54
            }
 
55
            __syncthreads();
 
56
        }
 
57
 
 
58
        if( nthreads == 512 )
 
59
            do_apply< 512 >( s );
 
60
        else if( nthreads == 256 )
 
61
            do_apply< 256 >( s );
 
62
        else if( nthreads == 128 )
 
63
            do_apply< 128 >( s );
 
64
        else if( nthreads == 64 )
 
65
            do_apply< 64 >( s );
 
66
        else if( nthreads == 32 )
 
67
            do_apply< 32 >( s );
 
68
        else if( nthreads == 16 )
 
69
            do_apply< 16 >( s );
 
70
        else if( nthreads == 8 )
 
71
            do_apply< 8 >( s );
 
72
        else if( nthreads == 4 )
 
73
            do_apply< 4 >( s );
 
74
        else if( nthreads == 2 )
 
75
            do_apply< 2 >( s );
 
76
    }
 
77
 
 
78
private:
 
79
   // This routine actually performs the reduction in shared memory
 
80
   // Result of the reduction is stored in s[0]
 
81
   template< unsigned int threads >
 
82
   static inline __device__ void do_apply( volatile float *s )
 
83
   {
 
84
      const unsigned int i = threadIdx.x;
 
85
 
 
86
      if( threads > 512 ) { if( i + 512 < threads ) s[i] = Op::eval( s[i], s[i + 512] ); __syncthreads(); }
 
87
      if( threads > 256 ) { if( i + 256 < threads ) s[i] = Op::eval( s[i], s[i + 256] ); __syncthreads(); }
 
88
      if( threads > 128 ) { if( i + 128 < threads ) s[i] = Op::eval( s[i], s[i + 128] ); __syncthreads(); }
 
89
      if( threads >  64 ) { if( i +  64 < threads ) s[i] = Op::eval( s[i], s[i +  64] ); __syncthreads(); }
 
90
 
 
91
      if( i < 32 )
 
92
      {
 
93
         // all this happens in one warp: no synchronization needed as s is declared volatile
 
94
         if( threads > 32 ) if( i + 32 < threads ) s[i] = Op::eval( s[i], s[i + 32] );
 
95
         if( threads > 16 ) if( i + 16 < threads ) s[i] = Op::eval( s[i], s[i + 16] );
 
96
         if( threads >  8 ) if( i +  8 < threads ) s[i] = Op::eval( s[i], s[i +  8] );
 
97
         if( threads >  4 ) if( i +  4 < threads ) s[i] = Op::eval( s[i], s[i +  4] );
 
98
         if( threads >  2 ) if( i +  2 < threads ) s[i] = Op::eval( s[i], s[i +  2] );
 
99
         if( threads >  1 ) if( i +  1 < threads ) s[i] = Op::eval( s[i], s[i +  1] );
 
100
      }
 
101
   }
 
102
};
 
103
 
 
104
#endif  // WCUDAREDUCE_CU