Commit 8dbe1eb4 authored by Tomáš Oberhuber's avatar Tomáš Oberhuber
Browse files

Fixing CUDA reduction kernels for the CUDA 3.0. 

Fixining yywrap() in debug scanner.
parent b4d79182

src/core/tnl-cuda-kernels.h

+12 −6
Original line number Diff line number Diff line
@@ -177,7 +177,8 @@ __global__ void tnlCUDAReductionKernel( const int size, 
                                        T* deviceOutput,

                                        T* dbg_array1 = 0 )

{

   extern __shared__ __align__ ( 8 ) T sdata[];

   extern __shared__ __align__ ( 8 ) char __sdata[];

   T* sdata = reinterpret_cast< T* >( __sdata );



   /***

    * Get thread id (tid) and global thread id (gid).
@@ -513,7 +514,8 @@ __global__ void tnlCUDASimpleReductionKernel5( const int size, 
                                               T* deviceOutput,

                                               T* dbg_array1 = 0 )

{

   extern __shared__ T sdata[];

   extern __shared__ __align__ ( 8 ) char __sdata[];

   T* sdata = reinterpret_cast< T* >( __sdata );



   /***

    * Get thread id (tid) and global thread id (gid).
@@ -732,7 +734,8 @@ __global__ void tnlCUDASimpleReductionKernel4( const int size, 
	                                       T* deviceOutput,

	                   	               T* dbg_array1 = 0  )

{

   extern __shared__ T sdata[];

   extern __shared__ __align__ ( 8 ) char __sdata[];

   T* sdata = reinterpret_cast< T* >( __sdata );



   /***

    * Get thread id (tid) and global thread id (gid).
@@ -965,7 +968,8 @@ __global__ void tnlCUDASimpleReductionKernel3( const int size, 
                                               const T* deviceInput,

		                               T* deviceOutput )

{

   extern __shared__ T sdata[];

   extern __shared__ __align__ ( 8 ) char __sdata[];

   T* sdata = reinterpret_cast< T* >( __sdata );



   /***

    * Get thread id (tid) and global thread id (gid).
@@ -1157,7 +1161,8 @@ __global__ void tnlCUDASimpleReductionKernel2( const int size, 
		                               const T* deviceInput,

		                               T* deviceOutput )

{

   extern __shared__ T sdata[];

   extern __shared__ __align__ ( 8 ) char __sdata[];

   T* sdata = reinterpret_cast< T* >( __sdata );



   /***

    * Get thread id (tid) and global thread id (gid).
@@ -1335,7 +1340,8 @@ __global__ void tnlCUDASimpleReductionKernel1( const int size, 
		                               const T* deviceInput,

		                               T* deviceOutput )

{

   extern __shared__ T sdata[];

   extern __shared__ __align__ ( 8 ) char __sdata[];

   T* sdata = reinterpret_cast< T* >( __sdata );



   /***

    * Get thread id (tid) and global thread id (gid).

src/debug/parse.cc

+1 −1
Original line number Diff line number Diff line 
// Generated by Bisonc++ V2.4.7 on Sat, 19 Dec 2009 18:04:01 +0100

// Generated by Bisonc++ V2.4.7 on Thu, 27 May 2010 13:23:09 +0200



// $insert class.ih

#include "tnlDebugParser.ih"

src/debug/tnlDebugParserbase.h

+1 −1
Original line number Diff line number Diff line 
// Generated by Bisonc++ V2.4.7 on Sat, 19 Dec 2009 18:04:01 +0100

// Generated by Bisonc++ V2.4.7 on Thu, 27 May 2010 13:23:09 +0200



#ifndef tnlDebugParserBase_h_included

#define tnlDebugParserBase_h_included

src/debug/tnlDebugScanner.cpp

+4 −4
Original line number Diff line number Diff line
@@ -1944,13 +1944,13 @@ void tnlDebugfree (void * ptr ) 






int tnlDebugFlexLexer::yywrap()

/*int tnlDebugFlexLexer::yywrap()

{

   return 1;

}

}*/



/*int yywrap()

{

extern int yywrap();

/*{

   return 1;

}*/

src/debug/tnlDebugScanner.l

+4 −4
Original line number Diff line number Diff line
@@ -41,12 +41,12 @@ default-debug { //cout << "DEFAULT" << endl; 
.                                             { return yytext[ 0 ]; }

%%



int tnlDebugFlexLexer::yywrap()

/*int tnlDebugFlexLexer::yywrap()

{

   return 1;

}

}*/



/*int yywrap()

{

extern int yywrap();

/*{

   return 1;

}*/