summaryrefslogtreecommitdiff
path: root/benchmarks/CUDA/RAY
diff options
context:
space:
mode:
Diffstat (limited to 'benchmarks/CUDA/RAY')
-rw-r--r--benchmarks/CUDA/RAY/EasyBMP.cpp1905
-rw-r--r--benchmarks/CUDA/RAY/EasyBMP.h86
-rw-r--r--benchmarks/CUDA/RAY/EasyBMP_BMP.h86
-rw-r--r--benchmarks/CUDA/RAY/EasyBMP_DataStructures.h104
-rw-r--r--benchmarks/CUDA/RAY/EasyBMP_VariousBMPutilities.h43
-rw-r--r--benchmarks/CUDA/RAY/Makefile50
-rw-r--r--benchmarks/CUDA/RAY/README.GPGPU-Sim2
-rw-r--r--benchmarks/CUDA/RAY/makebmp.cpp28
-rw-r--r--benchmarks/CUDA/RAY/makebmp.h9
-rw-r--r--benchmarks/CUDA/RAY/rayTracing.cu441
-rw-r--r--benchmarks/CUDA/RAY/rayTracing_kernel.cu658
11 files changed, 3412 insertions, 0 deletions
diff --git a/benchmarks/CUDA/RAY/EasyBMP.cpp b/benchmarks/CUDA/RAY/EasyBMP.cpp
new file mode 100644
index 0000000..c00122b
--- /dev/null
+++ b/benchmarks/CUDA/RAY/EasyBMP.cpp
@@ -0,0 +1,1905 @@
+/*************************************************
+* *
+* EasyBMP Cross-Platform Windows Bitmap Library *
+* *
+* Author: Paul Macklin *
+* email: [email protected] *
+* support: http://easybmp.sourceforge.net *
+* *
+* file: EasyBMP.cpp *
+* date added: 03-31-2006 *
+* date modified: 12-01-2006 *
+* version: 1.06 *
+* *
+* License: BSD (revised/modified) *
+* Copyright: 2005-6 by the EasyBMP Project *
+* *
+* description: Actual source file *
+* *
+*************************************************/
+
+#include "EasyBMP.h"
+
+/* These functions are defined in EasyBMP.h */
+
+bool EasyBMPwarnings = true;
+
+void SetEasyBMPwarningsOff( void )
+{ EasyBMPwarnings = false; }
+void SetEasyBMPwarningsOn( void )
+{ EasyBMPwarnings = true; }
+bool GetEasyBMPwarningState( void )
+{ return EasyBMPwarnings; }
+
+/* These functions are defined in EasyBMP_DataStructures.h */
+
+int IntPow( int base, int exponent )
+{
+ int i;
+ int output = 1;
+ for( i=0 ; i < exponent ; i++ )
+ { output *= base; }
+ return output;
+}
+
+BMFH::BMFH()
+{
+ bfType = 19778;
+ bfReserved1 = 0;
+ bfReserved2 = 0;
+}
+
+void BMFH::SwitchEndianess( void )
+{
+ bfType = FlipWORD( bfType );
+ bfSize = FlipDWORD( bfSize );
+ bfReserved1 = FlipWORD( bfReserved1 );
+ bfReserved2 = FlipWORD( bfReserved2 );
+ bfOffBits = FlipDWORD( bfOffBits );
+ return;
+}
+
+BMIH::BMIH()
+{
+ biPlanes = 1;
+ biCompression = 0;
+ biXPelsPerMeter = DefaultXPelsPerMeter;
+ biYPelsPerMeter = DefaultYPelsPerMeter;
+ biClrUsed = 0;
+ biClrImportant = 0;
+}
+
+void BMIH::SwitchEndianess( void )
+{
+ biSize = FlipDWORD( biSize );
+ biWidth = FlipDWORD( biWidth );
+ biHeight = FlipDWORD( biHeight );
+ biPlanes = FlipWORD( biPlanes );
+ biBitCount = FlipWORD( biBitCount );
+ biCompression = FlipDWORD( biCompression );
+ biSizeImage = FlipDWORD( biSizeImage );
+ biXPelsPerMeter = FlipDWORD( biXPelsPerMeter );
+ biYPelsPerMeter = FlipDWORD( biYPelsPerMeter );
+ biClrUsed = FlipDWORD( biClrUsed );
+ biClrImportant = FlipDWORD( biClrImportant );
+ return;
+}
+
+void BMIH::display( void )
+{
+ using namespace std;
+ cout << "biSize: " << (int) biSize << endl
+ << "biWidth: " << (int) biWidth << endl
+ << "biHeight: " << (int) biHeight << endl
+ << "biPlanes: " << (int) biPlanes << endl
+ << "biBitCount: " << (int) biBitCount << endl
+ << "biCompression: " << (int) biCompression << endl
+ << "biSizeImage: " << (int) biSizeImage << endl
+ << "biXPelsPerMeter: " << (int) biXPelsPerMeter << endl
+ << "biYPelsPerMeter: " << (int) biYPelsPerMeter << endl
+ << "biClrUsed: " << (int) biClrUsed << endl
+ << "biClrImportant: " << (int) biClrImportant << endl << endl;
+}
+
+void BMFH::display( void )
+{
+ using namespace std;
+ cout << "bfType: " << (int) bfType << endl
+ << "bfSize: " << (int) bfSize << endl
+ << "bfReserved1: " << (int) bfReserved1 << endl
+ << "bfReserved2: " << (int) bfReserved2 << endl
+ << "bfOffBits: " << (int) bfOffBits << endl << endl;
+}
+
+/* These functions are defined in EasyBMP_BMP.h */
+
+RGBApixel BMP::GetPixel( int i, int j ) const
+{
+ using namespace std;
+ bool Warn = false;
+ if( i >= Width )
+ { i = Width-1; Warn = true; }
+ if( i < 0 )
+ { i = 0; Warn = true; }
+ if( j >= Height )
+ { j = Height-1; Warn = true; }
+ if( j < 0 )
+ { j = 0; Warn = true; }
+ if( Warn && EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Attempted to access non-existent pixel;" << endl
+ << " Truncating request to fit in the range [0,"
+ << Width-1 << "] x [0," << Height-1 << "]." << endl;
+ }
+ return Pixels[i][j];
+}
+
+bool BMP::SetPixel( int i, int j, RGBApixel NewPixel )
+{
+ Pixels[i][j] = NewPixel;
+ return true;
+}
+
+
+bool BMP::SetColor( int ColorNumber , RGBApixel NewColor )
+{
+ using namespace std;
+ if( BitDepth != 1 && BitDepth != 4 && BitDepth != 8 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Attempted to change color table for a BMP object" << endl
+ << " that lacks a color table. Ignoring request." << endl;
+ }
+ return false;
+ }
+ if( !Colors )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Attempted to set a color, but the color table" << endl
+ << " is not defined. Ignoring request." << endl;
+ }
+ return false;
+ }
+ if( ColorNumber >= TellNumberOfColors() )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Requested color number "
+ << ColorNumber << " is outside the allowed" << endl
+ << " range [0," << TellNumberOfColors()-1
+ << "]. Ignoring request to set this color." << endl;
+ }
+ return false;
+ }
+ Colors[ColorNumber] = NewColor;
+ return true;
+}
+
+// RGBApixel BMP::GetColor( int ColorNumber ) const
+RGBApixel BMP::GetColor( int ColorNumber )
+{
+ RGBApixel Output;
+ Output.Red = 255;
+ Output.Green = 255;
+ Output.Blue = 255;
+ Output.Alpha = 0;
+
+ using namespace std;
+ if( BitDepth != 1 && BitDepth != 4 && BitDepth != 8 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Attempted to access color table for a BMP object" << endl
+ << " that lacks a color table. Ignoring request." << endl;
+ }
+ return Output;
+ }
+ if( !Colors )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Requested a color, but the color table" << endl
+ << " is not defined. Ignoring request." << endl;
+ }
+ return Output;
+ }
+ if( ColorNumber >= TellNumberOfColors() )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Requested color number "
+ << ColorNumber << " is outside the allowed" << endl
+ << " range [0," << TellNumberOfColors()-1
+ << "]. Ignoring request to get this color." << endl;
+ }
+ return Output;
+ }
+ Output = Colors[ColorNumber];
+ return Output;
+}
+
+BMP::BMP()
+{
+ Width = 1;
+ Height = 1;
+ BitDepth = 24;
+ Pixels = new RGBApixel* [Width];
+ Pixels[0] = new RGBApixel [Height];
+ Colors = NULL;
+
+ XPelsPerMeter = 0;
+ YPelsPerMeter = 0;
+
+ MetaData1 = NULL;
+ SizeOfMetaData1 = 0;
+ MetaData2 = NULL;
+ SizeOfMetaData2 = 0;
+}
+
+// BMP::BMP( const BMP& Input )
+BMP::BMP( BMP& Input )
+{
+ // first, make the image empty.
+
+ Width = 1;
+ Height = 1;
+ BitDepth = 24;
+ Pixels = new RGBApixel* [Width];
+ Pixels[0] = new RGBApixel [Height];
+ Colors = NULL;
+ XPelsPerMeter = 0;
+ YPelsPerMeter = 0;
+
+ MetaData1 = NULL;
+ SizeOfMetaData1 = 0;
+ MetaData2 = NULL;
+ SizeOfMetaData2 = 0;
+
+ // now, set the correct bit depth
+
+ SetBitDepth( Input.TellBitDepth() );
+
+ // set the correct pixel size
+
+ SetSize( Input.TellWidth() , Input.TellHeight() );
+
+ // set the DPI information from Input
+
+ SetDPI( Input.TellHorizontalDPI() , Input.TellVerticalDPI() );
+
+ // if there is a color table, get all the colors
+
+ if( BitDepth == 1 || BitDepth == 4 ||
+ BitDepth == 8 )
+ {
+ for( int k=0 ; k < TellNumberOfColors() ; k++ )
+ {
+ SetColor( k, Input.GetColor( k ));
+ }
+ }
+
+ // get all the pixels
+
+ for( int j=0; j < Height ; j++ )
+ {
+ for( int i=0; i < Width ; i++ )
+ {
+ Pixels[i][j] = *Input(i,j);
+// Pixels[i][j] = Input.GetPixel(i,j); // *Input(i,j);
+ }
+ }
+}
+
+BMP::~BMP()
+{
+ int i;
+ for(i=0;i<Width;i++)
+ { delete [] Pixels[i]; }
+ delete [] Pixels;
+ if( Colors )
+ { delete [] Colors; }
+
+ if( MetaData1 )
+ { delete [] MetaData1; }
+ if( MetaData2 )
+ { delete [] MetaData2; }
+}
+
+RGBApixel* BMP::operator()(int i, int j)
+{
+ using namespace std;
+ bool Warn = false;
+ if( i >= Width )
+ { i = Width-1; Warn = true; }
+ if( i < 0 )
+ { i = 0; Warn = true; }
+ if( j >= Height )
+ { j = Height-1; Warn = true; }
+ if( j < 0 )
+ { j = 0; Warn = true; }
+ if( Warn && EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Attempted to access non-existent pixel;" << endl
+ << " Truncating request to fit in the range [0,"
+ << Width-1 << "] x [0," << Height-1 << "]." << endl;
+ }
+ return &(Pixels[i][j]);
+}
+
+// int BMP::TellBitDepth( void ) const
+int BMP::TellBitDepth( void )
+{ return BitDepth; }
+
+// int BMP::TellHeight( void ) const
+int BMP::TellHeight( void )
+{ return Height; }
+
+// int BMP::TellWidth( void ) const
+int BMP::TellWidth( void )
+{ return Width; }
+
+// int BMP::TellNumberOfColors( void ) const
+int BMP::TellNumberOfColors( void )
+{
+ int output = IntPow( 2, BitDepth );
+ if( BitDepth == 32 )
+ { output = IntPow( 2, 24 ); }
+ return output;
+}
+
+bool BMP::SetBitDepth( int NewDepth )
+{
+ using namespace std;
+ if( NewDepth != 1 && NewDepth != 4 &&
+ NewDepth != 8 && NewDepth != 16 &&
+ NewDepth != 24 && NewDepth != 32 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: User attempted to set unsupported bit depth "
+ << NewDepth << "." << endl
+ << " Bit depth remains unchanged at "
+ << BitDepth << "." << endl;
+ }
+ return false;
+ }
+
+ BitDepth = NewDepth;
+ if( Colors )
+ { delete [] Colors; }
+ int NumberOfColors = IntPow( 2, BitDepth );
+ if( BitDepth == 1 || BitDepth == 4 || BitDepth == 8 )
+ { Colors = new RGBApixel [NumberOfColors]; }
+ else
+ { Colors = NULL; }
+ if( BitDepth == 1 || BitDepth == 4 || BitDepth == 8 )
+ { CreateStandardColorTable(); }
+
+ return true;
+}
+
+bool BMP::SetSize(int NewWidth , int NewHeight )
+{
+ using namespace std;
+ if( NewWidth <= 0 || NewHeight <= 0 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: User attempted to set a non-positive width or height." << endl
+ << " Size remains unchanged at "
+ << Width << " x " << Height << "." << endl;
+ }
+ return false;
+ }
+
+ int i,j;
+
+ for(i=0;i<Width;i++)
+ { delete [] Pixels[i]; }
+ delete [] Pixels;
+
+ Width = NewWidth;
+ Height = NewHeight;
+ Pixels = new RGBApixel* [ Width ];
+
+ for(i=0; i<Width; i++)
+ { Pixels[i] = new RGBApixel [ Height ]; }
+
+ for( i=0 ; i < Width ; i++)
+ {
+ for( j=0 ; j < Height ; j++ )
+ {
+ Pixels[i][j].Red = 255;
+ Pixels[i][j].Green = 255;
+ Pixels[i][j].Blue = 255;
+ Pixels[i][j].Alpha = 0;
+ }
+ }
+
+ return true;
+}
+
+bool BMP::WriteToFile( const char* FileName )
+{
+ using namespace std;
+ if( !EasyBMPcheckDataSize() )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Data types are wrong size!" << endl
+ << " You may need to mess with EasyBMP_DataTypes.h" << endl
+ << " to fix these errors, and then recompile." << endl
+ << " All 32-bit and 64-bit machines should be" << endl
+ << " supported, however." << endl << endl;
+ }
+ return false;
+ }
+
+ FILE* fp = fopen( FileName, "wb" );
+ if( fp == NULL )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Cannot open file "
+ << FileName << " for output." << endl;
+ }
+ fclose( fp );
+ return false;
+ }
+
+ // some preliminaries
+
+ double dBytesPerPixel = ( (double) BitDepth ) / 8.0;
+ double dBytesPerRow = dBytesPerPixel * (Width+0.0);
+ dBytesPerRow = ceil(dBytesPerRow);
+
+ int BytePaddingPerRow = 4 - ( (int) (dBytesPerRow) )% 4;
+ if( BytePaddingPerRow == 4 )
+ { BytePaddingPerRow = 0; }
+
+ double dActualBytesPerRow = dBytesPerRow + BytePaddingPerRow;
+
+ double dTotalPixelBytes = Height * dActualBytesPerRow;
+
+ double dPaletteSize = 0;
+ if( BitDepth == 1 || BitDepth == 4 || BitDepth == 8 )
+ { dPaletteSize = IntPow(2,BitDepth)*4.0; }
+
+ // leave some room for 16-bit masks
+ if( BitDepth == 16 )
+ { dPaletteSize = 3*4; }
+
+ double dTotalFileSize = 14 + 40 + dPaletteSize + dTotalPixelBytes;
+
+ // write the file header
+
+ BMFH bmfh;
+ bmfh.bfSize = (ebmpDWORD) dTotalFileSize;
+ bmfh.bfReserved1 = 0;
+ bmfh.bfReserved2 = 0;
+ bmfh.bfOffBits = (ebmpDWORD) (14+40+dPaletteSize);
+
+ if( IsBigEndian() )
+ { bmfh.SwitchEndianess(); }
+
+ fwrite( (char*) &(bmfh.bfType) , sizeof(ebmpWORD) , 1 , fp );
+ fwrite( (char*) &(bmfh.bfSize) , sizeof(ebmpDWORD) , 1 , fp );
+ fwrite( (char*) &(bmfh.bfReserved1) , sizeof(ebmpWORD) , 1 , fp );
+ fwrite( (char*) &(bmfh.bfReserved2) , sizeof(ebmpWORD) , 1 , fp );
+ fwrite( (char*) &(bmfh.bfOffBits) , sizeof(ebmpDWORD) , 1 , fp );
+
+ // write the info header
+
+ BMIH bmih;
+ bmih.biSize = 40;
+ bmih.biWidth = Width;
+ bmih.biHeight = Height;
+ bmih.biPlanes = 1;
+ bmih.biBitCount = BitDepth;
+ bmih.biCompression = 0;
+ bmih.biSizeImage = (ebmpDWORD) dTotalPixelBytes;
+ if( XPelsPerMeter )
+ { bmih.biXPelsPerMeter = XPelsPerMeter; }
+ else
+ { bmih.biXPelsPerMeter = DefaultXPelsPerMeter; }
+ if( YPelsPerMeter )
+ { bmih.biYPelsPerMeter = YPelsPerMeter; }
+ else
+ { bmih.biYPelsPerMeter = DefaultYPelsPerMeter; }
+
+ bmih.biClrUsed = 0;
+ bmih.biClrImportant = 0;
+
+ // indicates that we'll be using bit fields for 16-bit files
+ if( BitDepth == 16 )
+ { bmih.biCompression = 3; }
+
+ if( IsBigEndian() )
+ { bmih.SwitchEndianess(); }
+
+ fwrite( (char*) &(bmih.biSize) , sizeof(ebmpDWORD) , 1 , fp );
+ fwrite( (char*) &(bmih.biWidth) , sizeof(ebmpDWORD) , 1 , fp );
+ fwrite( (char*) &(bmih.biHeight) , sizeof(ebmpDWORD) , 1 , fp );
+ fwrite( (char*) &(bmih.biPlanes) , sizeof(ebmpWORD) , 1 , fp );
+ fwrite( (char*) &(bmih.biBitCount) , sizeof(ebmpWORD) , 1 , fp );
+ fwrite( (char*) &(bmih.biCompression) , sizeof(ebmpDWORD) , 1 , fp );
+ fwrite( (char*) &(bmih.biSizeImage) , sizeof(ebmpDWORD) , 1 , fp );
+ fwrite( (char*) &(bmih.biXPelsPerMeter) , sizeof(ebmpDWORD) , 1 , fp );
+ fwrite( (char*) &(bmih.biYPelsPerMeter) , sizeof(ebmpDWORD) , 1 , fp );
+ fwrite( (char*) &(bmih.biClrUsed) , sizeof(ebmpDWORD) , 1 , fp);
+ fwrite( (char*) &(bmih.biClrImportant) , sizeof(ebmpDWORD) , 1 , fp);
+
+ // write the palette
+ if( BitDepth == 1 || BitDepth == 4 || BitDepth == 8 )
+ {
+ int NumberOfColors = IntPow(2,BitDepth);
+
+ // if there is no palette, create one
+ if( !Colors )
+ {
+ if( !Colors )
+ { Colors = new RGBApixel [NumberOfColors]; }
+ CreateStandardColorTable();
+ }
+
+ int n;
+ for( n=0 ; n < NumberOfColors ; n++ )
+ { fwrite( (char*) &(Colors[n]) , 4 , 1 , fp ); }
+ }
+
+ // write the pixels
+ int i,j;
+ if( BitDepth != 16 )
+ {
+ ebmpBYTE* Buffer;
+ int BufferSize = (int) ( (Width*BitDepth)/8.0 );
+ while( 8*BufferSize < Width*BitDepth )
+ { BufferSize++; }
+ while( BufferSize % 4 )
+ { BufferSize++; }
+
+ Buffer = new ebmpBYTE [BufferSize];
+ for( j=0 ; j < BufferSize; j++ )
+ { Buffer[j] = 0; }
+
+ j=Height-1;
+
+ while( j > -1 )
+ {
+ bool Success = false;
+ if( BitDepth == 32 )
+ { Success = Write32bitRow( Buffer, BufferSize, j ); }
+ if( BitDepth == 24 )
+ { Success = Write24bitRow( Buffer, BufferSize, j ); }
+ if( BitDepth == 8 )
+ { Success = Write8bitRow( Buffer, BufferSize, j ); }
+ if( BitDepth == 4 )
+ { Success = Write4bitRow( Buffer, BufferSize, j ); }
+ if( BitDepth == 1 )
+ { Success = Write1bitRow( Buffer, BufferSize, j ); }
+ if( Success )
+ {
+ int BytesWritten = (int) fwrite( (char*) Buffer, 1, BufferSize, fp );
+ if( BytesWritten != BufferSize )
+ { Success = false; }
+ }
+ if( !Success )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Could not write proper amount of data." << endl;
+ }
+ j = -1;
+ }
+ j--;
+ }
+
+ delete [] Buffer;
+ }
+
+ if( BitDepth == 16 )
+ {
+ // write the bit masks
+
+ ebmpWORD BlueMask = 31; // bits 12-16
+ ebmpWORD GreenMask = 2016; // bits 6-11
+ ebmpWORD RedMask = 63488; // bits 1-5
+ ebmpWORD ZeroWORD;
+
+ if( IsBigEndian() )
+ { RedMask = FlipWORD( RedMask ); }
+ fwrite( (char*) &RedMask , 2 , 1 , fp );
+ fwrite( (char*) &ZeroWORD , 2 , 1 , fp );
+
+ if( IsBigEndian() )
+ { GreenMask = FlipWORD( GreenMask ); }
+ fwrite( (char*) &GreenMask , 2 , 1 , fp );
+ fwrite( (char*) &ZeroWORD , 2 , 1 , fp );
+
+ if( IsBigEndian() )
+ { BlueMask = FlipWORD( BlueMask ); }
+ fwrite( (char*) &BlueMask , 2 , 1 , fp );
+ fwrite( (char*) &ZeroWORD , 2 , 1 , fp );
+
+ int DataBytes = Width*2;
+ int PaddingBytes = ( 4 - DataBytes % 4 ) % 4;
+
+ // write the actual pixels
+
+ for( j=Height-1 ; j >= 0 ; j-- )
+ {
+ // write all row pixel data
+ i=0;
+ int WriteNumber = 0;
+ while( WriteNumber < DataBytes )
+ {
+ ebmpWORD TempWORD;
+
+ ebmpWORD RedWORD = (ebmpWORD) ((Pixels[i][j]).Red / 8);
+ ebmpWORD GreenWORD = (ebmpWORD) ((Pixels[i][j]).Green / 4);
+ ebmpWORD BlueWORD = (ebmpWORD) ((Pixels[i][j]).Blue / 8);
+
+ TempWORD = (RedWORD<<11) + (GreenWORD<<5) + BlueWORD;
+ if( IsBigEndian() )
+ { TempWORD = FlipWORD( TempWORD ); }
+
+ fwrite( (char*) &TempWORD , 2, 1, fp);
+ WriteNumber += 2;
+ i++;
+ }
+ // write any necessary row padding
+ WriteNumber = 0;
+ while( WriteNumber < PaddingBytes )
+ {
+ ebmpBYTE TempBYTE;
+ fwrite( (char*) &TempBYTE , 1, 1, fp);
+ WriteNumber++;
+ }
+ }
+
+ }
+
+ fclose(fp);
+ return true;
+}
+
+bool BMP::ReadFromFile( const char* FileName )
+{
+ using namespace std;
+ if( !EasyBMPcheckDataSize() )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Data types are wrong size!" << endl
+ << " You may need to mess with EasyBMP_DataTypes.h" << endl
+ << " to fix these errors, and then recompile." << endl
+ << " All 32-bit and 64-bit machines should be" << endl
+ << " supported, however." << endl << endl;
+ }
+ return false;
+ }
+
+ FILE* fp = fopen( FileName, "rb" );
+ if( fp == NULL )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Cannot open file "
+ << FileName << " for input." << endl;
+ }
+ SetBitDepth(1);
+ SetSize(1,1);
+ return false;
+ }
+
+ // read the file header
+
+ BMFH bmfh;
+ bool NotCorrupted = true;
+
+ NotCorrupted &= SafeFread( (char*) &(bmfh.bfType) , sizeof(ebmpWORD), 1, fp);
+
+ bool IsBitmap = false;
+
+ if( IsBigEndian() && bmfh.bfType == 16973 )
+ { IsBitmap = true; }
+ if( !IsBigEndian() && bmfh.bfType == 19778 )
+ { IsBitmap = true; }
+
+ if( !IsBitmap )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: " << FileName
+ << " is not a Windows BMP file!" << endl;
+ }
+ fclose( fp );
+ return false;
+ }
+
+ NotCorrupted &= SafeFread( (char*) &(bmfh.bfSize) , sizeof(ebmpDWORD) , 1, fp);
+ NotCorrupted &= SafeFread( (char*) &(bmfh.bfReserved1) , sizeof(ebmpWORD) , 1, fp);
+ NotCorrupted &= SafeFread( (char*) &(bmfh.bfReserved2) , sizeof(ebmpWORD) , 1, fp);
+ NotCorrupted &= SafeFread( (char*) &(bmfh.bfOffBits) , sizeof(ebmpDWORD) , 1 , fp);
+
+ if( IsBigEndian() )
+ { bmfh.SwitchEndianess(); }
+
+ // read the info header
+
+ BMIH bmih;
+
+ NotCorrupted &= SafeFread( (char*) &(bmih.biSize) , sizeof(ebmpDWORD) , 1 , fp);
+ NotCorrupted &= SafeFread( (char*) &(bmih.biWidth) , sizeof(ebmpDWORD) , 1 , fp);
+ NotCorrupted &= SafeFread( (char*) &(bmih.biHeight) , sizeof(ebmpDWORD) , 1 , fp);
+ NotCorrupted &= SafeFread( (char*) &(bmih.biPlanes) , sizeof(ebmpWORD) , 1, fp);
+ NotCorrupted &= SafeFread( (char*) &(bmih.biBitCount) , sizeof(ebmpWORD) , 1, fp);
+
+ NotCorrupted &= SafeFread( (char*) &(bmih.biCompression) , sizeof(ebmpDWORD) , 1 , fp);
+ NotCorrupted &= SafeFread( (char*) &(bmih.biSizeImage) , sizeof(ebmpDWORD) , 1 , fp);
+ NotCorrupted &= SafeFread( (char*) &(bmih.biXPelsPerMeter) , sizeof(ebmpDWORD) , 1 , fp);
+ NotCorrupted &= SafeFread( (char*) &(bmih.biYPelsPerMeter) , sizeof(ebmpDWORD) , 1 , fp);
+ NotCorrupted &= SafeFread( (char*) &(bmih.biClrUsed) , sizeof(ebmpDWORD) , 1 , fp);
+ NotCorrupted &= SafeFread( (char*) &(bmih.biClrImportant) , sizeof(ebmpDWORD) , 1 , fp);
+
+ if( IsBigEndian() )
+ { bmih.SwitchEndianess(); }
+
+ // a safety catch: if any of the header information didn't read properly, abort
+ // future idea: check to see if at least most is self-consistent
+
+ if( !NotCorrupted )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: " << FileName
+ << " is obviously corrupted." << endl;
+ }
+ SetSize(1,1);
+ SetBitDepth(1);
+ fclose(fp);
+ return false;
+ }
+
+ XPelsPerMeter = bmih.biXPelsPerMeter;
+ YPelsPerMeter = bmih.biYPelsPerMeter;
+
+ // if bmih.biCompression 1 or 2, then the file is RLE compressed
+
+ if( bmih.biCompression == 1 || bmih.biCompression == 2 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: " << FileName << " is (RLE) compressed." << endl
+ << " EasyBMP does not support compression." << endl;
+ }
+ SetSize(1,1);
+ SetBitDepth(1);
+ fclose(fp);
+ return false;
+ }
+
+ // if bmih.biCompression > 3, then something strange is going on
+ // it's probably an OS2 bitmap file.
+
+ if( bmih.biCompression > 3 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: " << FileName << " is in an unsupported format."
+ << endl
+ << " (bmih.biCompression = "
+ << bmih.biCompression << ")" << endl
+ << " The file is probably an old OS2 bitmap or corrupted."
+ << endl;
+ }
+ SetSize(1,1);
+ SetBitDepth(1);
+ fclose(fp);
+ return false;
+ }
+
+ if( bmih.biCompression == 3 && bmih.biBitCount != 16 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: " << FileName
+ << " uses bit fields and is not a" << endl
+ << " 16-bit file. This is not supported." << endl;
+ }
+ SetSize(1,1);
+ SetBitDepth(1);
+ fclose(fp);
+ return false;
+ }
+
+ // set the bit depth
+
+ int TempBitDepth = (int) bmih.biBitCount;
+ if( TempBitDepth != 1 && TempBitDepth != 4
+ && TempBitDepth != 8 && TempBitDepth != 16
+ && TempBitDepth != 24 && TempBitDepth != 32 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: " << FileName << " has unrecognized bit depth." << endl;
+ }
+ SetSize(1,1);
+ SetBitDepth(1);
+ fclose(fp);
+ return false;
+ }
+ SetBitDepth( (int) bmih.biBitCount );
+
+ // set the size
+
+ if( (int) bmih.biWidth <= 0 || (int) bmih.biHeight <= 0 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: " << FileName
+ << " has a non-positive width or height." << endl;
+ }
+ SetSize(1,1);
+ SetBitDepth(1);
+ fclose(fp);
+ return false;
+ }
+ SetSize( (int) bmih.biWidth , (int) bmih.biHeight );
+
+ // some preliminaries
+
+ double dBytesPerPixel = ( (double) BitDepth ) / 8.0;
+ double dBytesPerRow = dBytesPerPixel * (Width+0.0);
+ dBytesPerRow = ceil(dBytesPerRow);
+
+ int BytePaddingPerRow = 4 - ( (int) (dBytesPerRow) )% 4;
+ if( BytePaddingPerRow == 4 )
+ { BytePaddingPerRow = 0; }
+
+ // if < 16 bits, read the palette
+
+ if( BitDepth < 16 )
+ {
+ // determine the number of colors specified in the
+ // color table
+
+ int NumberOfColorsToRead = ((int) bmfh.bfOffBits - 54 )/4;
+ if( NumberOfColorsToRead > IntPow(2,BitDepth) )
+ { NumberOfColorsToRead = IntPow(2,BitDepth); }
+
+ if( NumberOfColorsToRead < TellNumberOfColors() )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: file " << FileName << " has an underspecified" << endl
+ << " color table. The table will be padded with extra" << endl
+ << " white (255,255,255,0) entries." << endl;
+ }
+ }
+
+ int n;
+ for( n=0; n < NumberOfColorsToRead ; n++ )
+ {
+ SafeFread( (char*) &(Colors[n]) , 4 , 1 , fp);
+ }
+ for( n=NumberOfColorsToRead ; n < TellNumberOfColors() ; n++ )
+ {
+ RGBApixel WHITE;
+ WHITE.Red = 255;
+ WHITE.Green = 255;
+ WHITE.Blue = 255;
+ WHITE.Alpha = 0;
+ SetColor( n , WHITE );
+ }
+
+
+ }
+
+ // skip blank data if bfOffBits so indicates
+
+ int BytesToSkip = bmfh.bfOffBits - 54;;
+ if( BitDepth < 16 )
+ { BytesToSkip -= 4*IntPow(2,BitDepth); }
+ if( BitDepth == 16 && bmih.biCompression == 3 )
+ { BytesToSkip -= 3*4; }
+ if( BytesToSkip < 0 )
+ { BytesToSkip = 0; }
+ if( BytesToSkip > 0 && BitDepth != 16 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Extra meta data detected in file " << FileName << endl
+ << " Data will be skipped." << endl;
+ }
+ ebmpBYTE* TempSkipBYTE;
+ TempSkipBYTE = new ebmpBYTE [BytesToSkip];
+ SafeFread( (char*) TempSkipBYTE , BytesToSkip , 1 , fp);
+ delete [] TempSkipBYTE;
+ }
+
+ // This code reads 1, 4, 8, 24, and 32-bpp files
+ // with a more-efficient buffered technique.
+
+ int i,j;
+ if( BitDepth != 16 )
+ {
+ int BufferSize = (int) ( (Width*BitDepth) / 8.0 );
+ while( 8*BufferSize < Width*BitDepth )
+ { BufferSize++; }
+ while( BufferSize % 4 )
+ { BufferSize++; }
+ ebmpBYTE* Buffer;
+ Buffer = new ebmpBYTE [BufferSize];
+ j= Height-1;
+ while( j > -1 )
+ {
+ int BytesRead = (int) fread( (char*) Buffer, 1, BufferSize, fp );
+ if( BytesRead < BufferSize )
+ {
+ j = -1;
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Could not read proper amount of data." << endl;
+ }
+ }
+ else
+ {
+ bool Success = false;
+ if( BitDepth == 1 )
+ { Success = Read1bitRow( Buffer, BufferSize, j ); }
+ if( BitDepth == 4 )
+ { Success = Read4bitRow( Buffer, BufferSize, j ); }
+ if( BitDepth == 8 )
+ { Success = Read8bitRow( Buffer, BufferSize, j ); }
+ if( BitDepth == 24 )
+ { Success = Read24bitRow( Buffer, BufferSize, j ); }
+ if( BitDepth == 32 )
+ { Success = Read32bitRow( Buffer, BufferSize, j ); }
+ if( !Success )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Could not read enough pixel data!" << endl;
+ }
+ j = -1;
+ }
+ }
+ j--;
+ }
+ delete [] Buffer;
+ }
+
+ if( BitDepth == 16 )
+ {
+ int DataBytes = Width*2;
+ int PaddingBytes = ( 4 - DataBytes % 4 ) % 4;
+
+ // set the default mask
+
+ ebmpWORD BlueMask = 31; // bits 12-16
+ ebmpWORD GreenMask = 992; // bits 7-11
+ ebmpWORD RedMask = 31744; // bits 2-6
+
+ // read the bit fields, if necessary, to
+ // override the default 5-5-5 mask
+
+ if( bmih.biCompression != 0 )
+ {
+ // read the three bit masks
+
+ ebmpWORD TempMaskWORD;
+ ebmpWORD ZeroWORD;
+
+ SafeFread( (char*) &RedMask , 2 , 1 , fp );
+ if( IsBigEndian() )
+ { RedMask = FlipWORD(RedMask); }
+ SafeFread( (char*) &TempMaskWORD , 2, 1, fp );
+
+ SafeFread( (char*) &GreenMask , 2 , 1 , fp );
+ if( IsBigEndian() )
+ { GreenMask = FlipWORD(GreenMask); }
+ SafeFread( (char*) &TempMaskWORD , 2, 1, fp );
+
+ SafeFread( (char*) &BlueMask , 2 , 1 , fp );
+ if( IsBigEndian() )
+ { BlueMask = FlipWORD(BlueMask); }
+ SafeFread( (char*) &TempMaskWORD , 2, 1, fp );
+ }
+
+ // read and skip any meta data
+
+ if( BytesToSkip > 0 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Extra meta data detected in file "
+ << FileName << endl
+ << " Data will be skipped." << endl;
+ }
+ ebmpBYTE* TempSkipBYTE;
+ TempSkipBYTE = new ebmpBYTE [BytesToSkip];
+ SafeFread( (char*) TempSkipBYTE , BytesToSkip , 1 , fp);
+ delete [] TempSkipBYTE;
+ }
+
+ // determine the red, green and blue shifts
+
+ int GreenShift = 0;
+ ebmpWORD TempShiftWORD = GreenMask;
+ while( TempShiftWORD > 31 )
+ { TempShiftWORD = TempShiftWORD>>1; GreenShift++; }
+ int BlueShift = 0;
+ TempShiftWORD = BlueMask;
+ while( TempShiftWORD > 31 )
+ { TempShiftWORD = TempShiftWORD>>1; BlueShift++; }
+ int RedShift = 0;
+ TempShiftWORD = RedMask;
+ while( TempShiftWORD > 31 )
+ { TempShiftWORD = TempShiftWORD>>1; RedShift++; }
+
+ // read the actual pixels
+
+ for( j=Height-1 ; j >= 0 ; j-- )
+ {
+ i=0;
+ int ReadNumber = 0;
+ while( ReadNumber < DataBytes )
+ {
+ ebmpWORD TempWORD;
+ SafeFread( (char*) &TempWORD , 2 , 1 , fp );
+ if( IsBigEndian() )
+ { TempWORD = FlipWORD(TempWORD); }
+ ReadNumber += 2;
+
+ ebmpWORD Red = RedMask & TempWORD;
+ ebmpWORD Green = GreenMask & TempWORD;
+ ebmpWORD Blue = BlueMask & TempWORD;
+
+ ebmpBYTE BlueBYTE = (ebmpBYTE) 8*(Blue>>BlueShift);
+ ebmpBYTE GreenBYTE = (ebmpBYTE) 8*(Green>>GreenShift);
+ ebmpBYTE RedBYTE = (ebmpBYTE) 8*(Red>>RedShift);
+
+ (Pixels[i][j]).Red = RedBYTE;
+ (Pixels[i][j]).Green = GreenBYTE;
+ (Pixels[i][j]).Blue = BlueBYTE;
+
+ i++;
+ }
+ ReadNumber = 0;
+ while( ReadNumber < PaddingBytes )
+ {
+ ebmpBYTE TempBYTE;
+ SafeFread( (char*) &TempBYTE , 1, 1, fp);
+ ReadNumber++;
+ }
+ }
+
+ }
+
+ fclose(fp);
+ return true;
+}
+
+bool BMP::CreateStandardColorTable( void )
+{
+ using namespace std;
+ if( BitDepth != 1 && BitDepth != 4 && BitDepth != 8 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Attempted to create color table at a bit" << endl
+ << " depth that does not require a color table." << endl
+ << " Ignoring request." << endl;
+ }
+ return false;
+ }
+
+ if( BitDepth == 1 )
+ {
+ int i;
+ for( i=0 ; i < 2 ; i++ )
+ {
+ Colors[i].Red = i*255;
+ Colors[i].Green = i*255;
+ Colors[i].Blue = i*255;
+ Colors[i].Alpha = 0;
+ }
+ return true;
+ }
+
+ if( BitDepth == 4 )
+ {
+ int i = 0;
+ int j,k,ell;
+
+ // simplify the code for the first 8 colors
+ for( ell=0 ; ell < 2 ; ell++ )
+ {
+ for( k=0 ; k < 2 ; k++ )
+ {
+ for( j=0 ; j < 2 ; j++ )
+ {
+ Colors[i].Red = j*128;
+ Colors[i].Green = k*128;
+ Colors[i].Blue = ell*128;
+ i++;
+ }
+ }
+ }
+
+ // simplify the code for the last 8 colors
+ for( ell=0 ; ell < 2 ; ell++ )
+ {
+ for( k=0 ; k < 2 ; k++ )
+ {
+ for( j=0 ; j < 2 ; j++ )
+ {
+ Colors[i].Red = j*255;
+ Colors[i].Green = k*255;
+ Colors[i].Blue = ell*255;
+ i++;
+ }
+ }
+ }
+
+ // overwrite the duplicate color
+ i=8;
+ Colors[i].Red = 192;
+ Colors[i].Green = 192;
+ Colors[i].Blue = 192;
+
+ for( i=0 ; i < 16 ; i++ )
+ { Colors[i].Alpha = 0; }
+ return true;
+ }
+
+ if( BitDepth == 8 )
+ {
+ int i=0;
+ int j,k,ell;
+
+ // do an easy loop, which works for all but colors
+ // 0 to 9 and 246 to 255
+ for( ell=0 ; ell < 4 ; ell++ )
+ {
+ for( k=0 ; k < 8 ; k++ )
+ {
+ for( j=0; j < 8 ; j++ )
+ {
+ Colors[i].Red = j*32;
+ Colors[i].Green = k*32;
+ Colors[i].Blue = ell*64;
+ Colors[i].Alpha = 0;
+ i++;
+ }
+ }
+ }
+
+ // now redo the first 8 colors
+ i=0;
+ for( ell=0 ; ell < 2 ; ell++ )
+ {
+ for( k=0 ; k < 2 ; k++ )
+ {
+ for( j=0; j < 2 ; j++ )
+ {
+ Colors[i].Red = j*128;
+ Colors[i].Green = k*128;
+ Colors[i].Blue = ell*128;
+ i++;
+ }
+ }
+ }
+
+ // overwrite colors 7, 8, 9
+ i=7;
+ Colors[i].Red = 192;
+ Colors[i].Green = 192;
+ Colors[i].Blue = 192;
+ i++; // 8
+ Colors[i].Red = 192;
+ Colors[i].Green = 220;
+ Colors[i].Blue = 192;
+ i++; // 9
+ Colors[i].Red = 166;
+ Colors[i].Green = 202;
+ Colors[i].Blue = 240;
+
+ // overwrite colors 246 to 255
+ i=246;
+ Colors[i].Red = 255;
+ Colors[i].Green = 251;
+ Colors[i].Blue = 240;
+ i++; // 247
+ Colors[i].Red = 160;
+ Colors[i].Green = 160;
+ Colors[i].Blue = 164;
+ i++; // 248
+ Colors[i].Red = 128;
+ Colors[i].Green = 128;
+ Colors[i].Blue = 128;
+ i++; // 249
+ Colors[i].Red = 255;
+ Colors[i].Green = 0;
+ Colors[i].Blue = 0;
+ i++; // 250
+ Colors[i].Red = 0;
+ Colors[i].Green = 255;
+ Colors[i].Blue = 0;
+ i++; // 251
+ Colors[i].Red = 255;
+ Colors[i].Green = 255;
+ Colors[i].Blue = 0;
+ i++; // 252
+ Colors[i].Red = 0;
+ Colors[i].Green = 0;
+ Colors[i].Blue = 255;
+ i++; // 253
+ Colors[i].Red = 255;
+ Colors[i].Green = 0;
+ Colors[i].Blue = 255;
+ i++; // 254
+ Colors[i].Red = 0;
+ Colors[i].Green = 255;
+ Colors[i].Blue = 255;
+ i++; // 255
+ Colors[i].Red = 255;
+ Colors[i].Green = 255;
+ Colors[i].Blue = 255;
+
+ return true;
+ }
+ return true;
+}
+
+bool SafeFread( char* buffer, int size, int number, FILE* fp )
+{
+ using namespace std;
+ int ItemsRead;
+ if( feof(fp) )
+ { return false; }
+ ItemsRead = (int) fread( buffer , size , number , fp );
+ if( ItemsRead < number )
+ { return false; }
+ return true;
+}
+
+void BMP::SetDPI( int HorizontalDPI, int VerticalDPI )
+{
+ XPelsPerMeter = (int) ( HorizontalDPI * 39.37007874015748 );
+ YPelsPerMeter = (int) ( VerticalDPI * 39.37007874015748 );
+}
+
+// int BMP::TellVerticalDPI( void ) const
+int BMP::TellVerticalDPI( void )
+{
+ if( !YPelsPerMeter )
+ { YPelsPerMeter = DefaultYPelsPerMeter; }
+ return (int) ( YPelsPerMeter / (double) 39.37007874015748 );
+}
+
+// int BMP::TellHorizontalDPI( void ) const
+int BMP::TellHorizontalDPI( void )
+{
+ if( !XPelsPerMeter )
+ { XPelsPerMeter = DefaultXPelsPerMeter; }
+ return (int) ( XPelsPerMeter / (double) 39.37007874015748 );
+}
+
+/* These functions are defined in EasyBMP_VariousBMPutilities.h */
+
+BMFH GetBMFH( const char* szFileNameIn )
+{
+ using namespace std;
+ BMFH bmfh;
+
+ FILE* fp;
+ fp = fopen( szFileNameIn,"rb");
+
+ if( !fp )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Cannot initialize from file "
+ << szFileNameIn << "." << endl
+ << " File cannot be opened or does not exist."
+ << endl;
+ }
+ bmfh.bfType = 0;
+ return bmfh;
+ }
+
+ SafeFread( (char*) &(bmfh.bfType) , sizeof(ebmpWORD) , 1 , fp );
+ SafeFread( (char*) &(bmfh.bfSize) , sizeof(ebmpDWORD) , 1 , fp );
+ SafeFread( (char*) &(bmfh.bfReserved1) , sizeof(ebmpWORD) , 1 , fp );
+ SafeFread( (char*) &(bmfh.bfReserved2) , sizeof(ebmpWORD) , 1 , fp );
+ SafeFread( (char*) &(bmfh.bfOffBits) , sizeof(ebmpDWORD) , 1 , fp );
+
+ fclose( fp );
+
+ if( IsBigEndian() )
+ { bmfh.SwitchEndianess(); }
+
+ return bmfh;
+}
+
+BMIH GetBMIH( const char* szFileNameIn )
+{
+ using namespace std;
+ BMFH bmfh;
+ BMIH bmih;
+
+ FILE* fp;
+ fp = fopen( szFileNameIn,"rb");
+
+ if( !fp )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Cannot initialize from file "
+ << szFileNameIn << "." << endl
+ << " File cannot be opened or does not exist."
+ << endl;
+ }
+ return bmih;
+ }
+
+ // read the bmfh, i.e., first 14 bytes (just to get it out of the way);
+
+ ebmpBYTE TempBYTE;
+ int i;
+ for( i = 14 ; i > 0 ; i-- )
+ { SafeFread( (char*) &TempBYTE , sizeof(ebmpBYTE) , 1, fp ); }
+
+ // read the bmih
+
+ SafeFread( (char*) &(bmih.biSize) , sizeof(ebmpDWORD) , 1 , fp );
+ SafeFread( (char*) &(bmih.biWidth) , sizeof(ebmpDWORD) , 1 , fp );
+ SafeFread( (char*) &(bmih.biHeight) , sizeof(ebmpDWORD) , 1 , fp );
+ SafeFread( (char*) &(bmih.biPlanes) , sizeof(ebmpWORD) , 1 , fp );
+
+ SafeFread( (char*) &(bmih.biBitCount) , sizeof(ebmpWORD) , 1 , fp );
+ SafeFread( (char*) &(bmih.biCompression) , sizeof(ebmpDWORD) , 1 , fp );
+ SafeFread( (char*) &(bmih.biSizeImage) , sizeof(ebmpDWORD) , 1 , fp );
+ SafeFread( (char*) &(bmih.biXPelsPerMeter) , sizeof(ebmpDWORD) , 1 , fp );
+
+ SafeFread( (char*) &(bmih.biYPelsPerMeter) , sizeof(ebmpDWORD) , 1 , fp );
+ SafeFread( (char*) &(bmih.biClrUsed) , sizeof(ebmpDWORD) , 1 , fp );
+ SafeFread( (char*) &(bmih.biClrImportant) , sizeof(ebmpDWORD) , 1 , fp );
+
+ fclose( fp );
+
+ if( IsBigEndian() )
+ { bmih.SwitchEndianess(); }
+
+ return bmih;
+}
+
+void DisplayBitmapInfo( const char* szFileNameIn )
+{
+ using namespace std;
+ FILE* fp;
+ fp = fopen( szFileNameIn,"rb");
+
+ if( !fp )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: Cannot initialize from file "
+ << szFileNameIn << "." << endl
+ << " File cannot be opened or does not exist."
+ << endl;
+ }
+ return;
+ }
+ fclose( fp );
+
+ // don't duplicate work! Just use the functions from above!
+
+ BMFH bmfh = GetBMFH(szFileNameIn);
+ BMIH bmih = GetBMIH(szFileNameIn);
+
+ cout << "File information for file " << szFileNameIn
+ << ":" << endl << endl;
+
+ cout << "BITMAPFILEHEADER:" << endl
+ << "bfType: " << bmfh.bfType << endl
+ << "bfSize: " << bmfh.bfSize << endl
+ << "bfReserved1: " << bmfh.bfReserved1 << endl
+ << "bfReserved2: " << bmfh.bfReserved2 << endl
+ << "bfOffBits: " << bmfh.bfOffBits << endl << endl;
+
+ cout << "BITMAPINFOHEADER:" << endl
+ << "biSize: " << bmih.biSize << endl
+ << "biWidth: " << bmih.biWidth << endl
+ << "biHeight: " << bmih.biHeight << endl
+ << "biPlanes: " << bmih.biPlanes << endl
+ << "biBitCount: " << bmih.biBitCount << endl
+ << "biCompression: " << bmih.biCompression << endl
+ << "biSizeImage: " << bmih.biSizeImage << endl
+ << "biXPelsPerMeter: " << bmih.biXPelsPerMeter << endl
+ << "biYPelsPerMeter: " << bmih.biYPelsPerMeter << endl
+ << "biClrUsed: " << bmih.biClrUsed << endl
+ << "biClrImportant: " << bmih.biClrImportant << endl << endl;
+ return;
+}
+
+int GetBitmapColorDepth( const char* szFileNameIn )
+{
+ BMIH bmih = GetBMIH( szFileNameIn );
+ return (int) bmih.biBitCount;
+}
+
+void PixelToPixelCopy( BMP& From, int FromX, int FromY,
+ BMP& To, int ToX, int ToY)
+{
+ *To(ToX,ToY) = *From(FromX,FromY);
+ return;
+}
+
+void PixelToPixelCopyTransparent( BMP& From, int FromX, int FromY,
+ BMP& To, int ToX, int ToY,
+ RGBApixel& Transparent )
+{
+ if( From(FromX,FromY)->Red != Transparent.Red ||
+ From(FromX,FromY)->Green != Transparent.Green ||
+ From(FromX,FromY)->Blue != Transparent.Blue )
+ { *To(ToX,ToY) = *From(FromX,FromY); }
+ return;
+}
+
+void RangedPixelToPixelCopy( BMP& From, int FromL , int FromR, int FromB, int FromT,
+ BMP& To, int ToX, int ToY )
+{
+ // make sure the conventions are followed
+ if( FromB < FromT )
+ { int Temp = FromT; FromT = FromB; FromB = Temp; }
+
+ // make sure that the copied regions exist in both bitmaps
+ if( FromR >= From.TellWidth() )
+ { FromR = From.TellWidth()-1; }
+ if( FromL < 0 ){ FromL = 0; }
+
+ if( FromB >= From.TellHeight() )
+ { FromB = From.TellHeight()-1; }
+ if( FromT < 0 ){ FromT = 0; }
+
+ if( ToX+(FromR-FromL) >= To.TellWidth() )
+ { FromR = To.TellWidth()-1+FromL-ToX; }
+ if( ToY+(FromB-FromT) >= To.TellHeight() )
+ { FromB = To.TellHeight()-1+FromT-ToY; }
+
+ int i,j;
+ for( j=FromT ; j <= FromB ; j++ )
+ {
+ for( i=FromL ; i <= FromR ; i++ )
+ {
+ PixelToPixelCopy( From, i,j,
+ To, ToX+(i-FromL), ToY+(j-FromT) );
+ }
+ }
+
+ return;
+}
+
+void RangedPixelToPixelCopyTransparent(
+ BMP& From, int FromL , int FromR, int FromB, int FromT,
+ BMP& To, int ToX, int ToY ,
+ RGBApixel& Transparent )
+{
+ // make sure the conventions are followed
+ if( FromB < FromT )
+ { int Temp = FromT; FromT = FromB; FromB = Temp; }
+
+ // make sure that the copied regions exist in both bitmaps
+ if( FromR >= From.TellWidth() )
+ { FromR = From.TellWidth()-1; }
+ if( FromL < 0 ){ FromL = 0; }
+
+ if( FromB >= From.TellHeight() )
+ { FromB = From.TellHeight()-1; }
+ if( FromT < 0 ){ FromT = 0; }
+
+ if( ToX+(FromR-FromL) >= To.TellWidth() )
+ { FromR = To.TellWidth()-1+FromL-ToX; }
+ if( ToY+(FromB-FromT) >= To.TellHeight() )
+ { FromB = To.TellHeight()-1+FromT-ToY; }
+
+ int i,j;
+ for( j=FromT ; j <= FromB ; j++ )
+ {
+ for( i=FromL ; i <= FromR ; i++ )
+ {
+ PixelToPixelCopyTransparent( From, i,j,
+ To, ToX+(i-FromL), ToY+(j-FromT) ,
+ Transparent);
+ }
+ }
+
+ return;
+}
+
+bool CreateGrayscaleColorTable( BMP& InputImage )
+{
+ using namespace std;
+ int BitDepth = InputImage.TellBitDepth();
+ if( BitDepth != 1 && BitDepth != 4 && BitDepth != 8 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Warning: Attempted to create color table at a bit" << endl
+ << " depth that does not require a color table." << endl
+ << " Ignoring request." << endl;
+ }
+ return false;
+ }
+ int i;
+ int NumberOfColors = InputImage.TellNumberOfColors();
+
+ ebmpBYTE StepSize;
+ if( BitDepth != 1 )
+ { StepSize = 255/(NumberOfColors-1); }
+ else
+ { StepSize = 255; }
+
+ for( i=0 ; i < NumberOfColors ; i++ )
+ {
+ ebmpBYTE TempBYTE = i*StepSize;
+ RGBApixel TempColor;
+ TempColor.Red = TempBYTE;
+ TempColor.Green = TempBYTE;
+ TempColor.Blue = TempBYTE;
+ TempColor.Alpha = 0;
+ InputImage.SetColor( i , TempColor );
+ }
+ return true;
+}
+
+bool BMP::Read32bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int i;
+ if( Width*4 > BufferSize )
+ { return false; }
+ for( i=0 ; i < Width ; i++ )
+ { memcpy( (char*) &(Pixels[i][Row]), (char*) Buffer+4*i, 4 ); }
+ return true;
+}
+
+bool BMP::Read24bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int i;
+ if( Width*3 > BufferSize )
+ { return false; }
+ for( i=0 ; i < Width ; i++ )
+ { memcpy( (char*) &(Pixels[i][Row]), Buffer+3*i, 3 ); }
+ return true;
+}
+
+bool BMP::Read8bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int i;
+ if( Width > BufferSize )
+ { return false; }
+ for( i=0 ; i < Width ; i++ )
+ {
+ int Index = Buffer[i];
+ *( this->operator()(i,Row) )= GetColor(Index);
+ }
+ return true;
+}
+
+bool BMP::Read4bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int Shifts[2] = {4 ,0 };
+ int Masks[2] = {240,15};
+
+ int i=0;
+ int j;
+ int k=0;
+ if( Width > 2*BufferSize )
+ { return false; }
+ while( i < Width )
+ {
+ j=0;
+ while( j < 2 && i < Width )
+ {
+ int Index = (int) ( (Buffer[k]&Masks[j]) >> Shifts[j]);
+ *( this->operator()(i,Row) )= GetColor(Index);
+ i++; j++;
+ }
+ k++;
+ }
+ return true;
+}
+bool BMP::Read1bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int Shifts[8] = {7 ,6 ,5 ,4 ,3,2,1,0};
+ int Masks[8] = {128,64,32,16,8,4,2,1};
+
+ int i=0;
+ int j;
+ int k=0;
+
+ if( Width > 8*BufferSize )
+ { return false; }
+ while( i < Width )
+ {
+ j=0;
+ while( j < 8 && i < Width )
+ {
+ int Index = (int) ( (Buffer[k]&Masks[j]) >> Shifts[j]);
+ *( this->operator()(i,Row) )= GetColor(Index);
+ i++; j++;
+ }
+ k++;
+ }
+ return true;
+}
+
+bool BMP::Write32bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int i;
+ if( Width*4 > BufferSize )
+ { return false; }
+ for( i=0 ; i < Width ; i++ )
+ { memcpy( (char*) Buffer+4*i, (char*) &(Pixels[i][Row]), 4 ); }
+ return true;
+}
+
+bool BMP::Write24bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int i;
+ if( Width*3 > BufferSize )
+ { return false; }
+ for( i=0 ; i < Width ; i++ )
+ { memcpy( (char*) Buffer+3*i, (char*) &(Pixels[i][Row]), 3 ); }
+ return true;
+}
+
+bool BMP::Write8bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int i;
+ if( Width > BufferSize )
+ { return false; }
+ for( i=0 ; i < Width ; i++ )
+ { Buffer[i] = FindClosestColor( Pixels[i][Row] ); }
+ return true;
+}
+
+bool BMP::Write4bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int PositionWeights[2] = {16,1};
+
+ int i=0;
+ int j;
+ int k=0;
+ if( Width > 2*BufferSize )
+ { return false; }
+ while( i < Width )
+ {
+ j=0;
+ int Index = 0;
+ while( j < 2 && i < Width )
+ {
+ Index += ( PositionWeights[j]* (int) FindClosestColor( Pixels[i][Row] ) );
+ i++; j++;
+ }
+ Buffer[k] = (ebmpBYTE) Index;
+ k++;
+ }
+ return true;
+}
+
+bool BMP::Write1bitRow( ebmpBYTE* Buffer, int BufferSize, int Row )
+{
+ int PositionWeights[8] = {128,64,32,16,8,4,2,1};
+
+ int i=0;
+ int j;
+ int k=0;
+ if( Width > 8*BufferSize )
+ { return false; }
+ while( i < Width )
+ {
+ j=0;
+ int Index = 0;
+ while( j < 8 && i < Width )
+ {
+ Index += ( PositionWeights[j]* (int) FindClosestColor( Pixels[i][Row] ) );
+ i++; j++;
+ }
+ Buffer[k] = (ebmpBYTE) Index;
+ k++;
+ }
+ return true;
+}
+
+ebmpBYTE BMP::FindClosestColor( RGBApixel& input )
+{
+ using namespace std;
+
+ int i=0;
+ int NumberOfColors = TellNumberOfColors();
+ ebmpBYTE BestI = 0;
+ int BestMatch = 999999;
+
+ while( i < NumberOfColors )
+ {
+ RGBApixel Attempt = GetColor( i );
+ int TempMatch = IntSquare( (int) Attempt.Red - (int) input.Red )
+ + IntSquare( (int) Attempt.Green - (int) input.Green )
+ + IntSquare( (int) Attempt.Blue - (int) input.Blue );
+ if( TempMatch < BestMatch )
+ { BestI = (ebmpBYTE) i; BestMatch = TempMatch; }
+ if( BestMatch < 1 )
+ { i = NumberOfColors; }
+ i++;
+ }
+ return BestI;
+}
+
+bool EasyBMPcheckDataSize( void )
+{
+ using namespace std;
+ bool ReturnValue = true;
+ if( sizeof( ebmpBYTE ) != 1 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: ebmpBYTE has the wrong size ("
+ << sizeof( ebmpBYTE ) << " bytes)," << endl
+ << " Compared to the expected 1 byte value" << endl;
+ }
+ ReturnValue = false;
+ }
+ if( sizeof( ebmpWORD ) != 2 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: ebmpWORD has the wrong size ("
+ << sizeof( ebmpWORD ) << " bytes)," << endl
+ << " Compared to the expected 2 byte value" << endl;
+ }
+ ReturnValue = false;
+ }
+ if( sizeof( ebmpDWORD ) != 4 )
+ {
+ if( EasyBMPwarnings )
+ {
+ cout << "EasyBMP Error: ebmpDWORD has the wrong size ("
+ << sizeof( ebmpDWORD ) << " bytes)," << endl
+ << " Compared to the expected 4 byte value" << endl;
+ }
+ ReturnValue = false;
+ }
+ return ReturnValue;
+}
+
+bool Rescale( BMP& InputImage , char mode, int NewDimension )
+{
+ using namespace std;
+ int CapMode = toupper( mode );
+
+ BMP OldImage( InputImage );
+
+ if( CapMode != 'P' &&
+ CapMode != 'W' &&
+ CapMode != 'H' &&
+ CapMode != 'F' )
+ {
+ if( EasyBMPwarnings )
+ {
+ char ErrorMessage [1024];
+ sprintf( ErrorMessage, "EasyBMP Error: Unknown rescale mode %c requested\n" , mode );
+ cout << ErrorMessage;
+ }
+ return false;
+ }
+
+ int NewWidth =0;
+ int NewHeight =0;
+
+ int OldWidth = OldImage.TellWidth();
+ int OldHeight= OldImage.TellHeight();
+
+ if( CapMode == 'P' )
+ {
+ NewWidth = (int) floor( OldWidth * NewDimension / 100.0 );
+ NewHeight = (int) floor( OldHeight * NewDimension / 100.0 );
+ }
+ if( CapMode == 'F' )
+ {
+ if( OldWidth > OldHeight )
+ { CapMode = 'W'; }
+ else
+ { CapMode = 'H'; }
+ }
+
+ if( CapMode == 'W' )
+ {
+ double percent = (double) NewDimension / (double) OldWidth;
+ NewWidth = NewDimension;
+ NewHeight = (int) floor( OldHeight * percent );
+ }
+ if( CapMode == 'H' )
+ {
+ double percent = (double) NewDimension / (double) OldHeight;
+ NewHeight = NewDimension;
+ NewWidth = (int) floor( OldWidth * percent );
+ }
+
+ if( NewWidth < 1 )
+ { NewWidth = 1; }
+ if( NewHeight < 1 )
+ { NewHeight = 1; }
+
+ InputImage.SetSize( NewWidth, NewHeight );
+ InputImage.SetBitDepth( 24 );
+
+ int I,J;
+ double ThetaI,ThetaJ;
+
+ for( int j=0; j < NewHeight-1 ; j++ )
+ {
+ ThetaJ = (double)(j*(OldHeight-1.0))
+ /(double)(NewHeight-1.0);
+ J = (int) floor( ThetaJ );
+ ThetaJ -= J;
+
+ for( int i=0; i < NewWidth-1 ; i++ )
+ {
+ ThetaI = (double)(i*(OldWidth-1.0))
+ /(double)(NewWidth-1.0);
+ I = (int) floor( ThetaI );
+ ThetaI -= I;
+
+ InputImage(i,j)->Red = (ebmpBYTE)
+ ( (1.0-ThetaI-ThetaJ+ThetaI*ThetaJ)*(OldImage(I,J)->Red)
+ +(ThetaI-ThetaI*ThetaJ)*(OldImage(I+1,J)->Red)
+ +(ThetaJ-ThetaI*ThetaJ)*(OldImage(I,J+1)->Red)
+ +(ThetaI*ThetaJ)*(OldImage(I+1,J+1)->Red) );
+ InputImage(i,j)->Green = (ebmpBYTE)
+ ( (1.0-ThetaI-ThetaJ+ThetaI*ThetaJ)*OldImage(I,J)->Green
+ +(ThetaI-ThetaI*ThetaJ)*OldImage(I+1,J)->Green
+ +(ThetaJ-ThetaI*ThetaJ)*OldImage(I,J+1)->Green
+ +(ThetaI*ThetaJ)*OldImage(I+1,J+1)->Green );
+ InputImage(i,j)->Blue = (ebmpBYTE)
+ ( (1.0-ThetaI-ThetaJ+ThetaI*ThetaJ)*OldImage(I,J)->Blue
+ +(ThetaI-ThetaI*ThetaJ)*OldImage(I+1,J)->Blue
+ +(ThetaJ-ThetaI*ThetaJ)*OldImage(I,J+1)->Blue
+ +(ThetaI*ThetaJ)*OldImage(I+1,J+1)->Blue );
+ }
+ InputImage(NewWidth-1,j)->Red = (ebmpBYTE)
+ ( (1.0-ThetaJ)*(OldImage(OldWidth-1,J)->Red)
+ + ThetaJ*(OldImage(OldWidth-1,J+1)->Red) );
+ InputImage(NewWidth-1,j)->Green = (ebmpBYTE)
+ ( (1.0-ThetaJ)*(OldImage(OldWidth-1,J)->Green)
+ + ThetaJ*(OldImage(OldWidth-1,J+1)->Green) );
+ InputImage(NewWidth-1,j)->Blue = (ebmpBYTE)
+ ( (1.0-ThetaJ)*(OldImage(OldWidth-1,J)->Blue)
+ + ThetaJ*(OldImage(OldWidth-1,J+1)->Blue) );
+ }
+
+ for( int i=0 ; i < NewWidth-1 ; i++ )
+ {
+ ThetaI = (double)(i*(OldWidth-1.0))
+ /(double)(NewWidth-1.0);
+ I = (int) floor( ThetaI );
+ ThetaI -= I;
+ InputImage(i,NewHeight-1)->Red = (ebmpBYTE)
+ ( (1.0-ThetaI)*(OldImage(I,OldHeight-1)->Red)
+ + ThetaI*(OldImage(I,OldHeight-1)->Red) );
+ InputImage(i,NewHeight-1)->Green = (ebmpBYTE)
+ ( (1.0-ThetaI)*(OldImage(I,OldHeight-1)->Green)
+ + ThetaI*(OldImage(I,OldHeight-1)->Green) );
+ InputImage(i,NewHeight-1)->Blue = (ebmpBYTE)
+ ( (1.0-ThetaI)*(OldImage(I,OldHeight-1)->Blue)
+ + ThetaI*(OldImage(I,OldHeight-1)->Blue) );
+ }
+
+ *InputImage(NewWidth-1,NewHeight-1) = *OldImage(OldWidth-1,OldHeight-1);
+ return true;
+}
diff --git a/benchmarks/CUDA/RAY/EasyBMP.h b/benchmarks/CUDA/RAY/EasyBMP.h
new file mode 100644
index 0000000..ead98c1
--- /dev/null
+++ b/benchmarks/CUDA/RAY/EasyBMP.h
@@ -0,0 +1,86 @@
+/*************************************************
+* *
+* EasyBMP Cross-Platform Windows Bitmap Library *
+* *
+* Author: Paul Macklin *
+* email: [email protected] *
+* support: http://easybmp.sourceforge.net *
+* *
+* file: EasyBMP.h *
+* date added: 01-31-2005 *
+* date modified: 12-01-2006 *
+* version: 1.06 *
+* *
+* License: BSD (revised/modified) *
+* Copyright: 2005-6 by the EasyBMP Project *
+* *
+* description: Main include file *
+* *
+*************************************************/
+
+#ifdef _MSC_VER
+// MS Visual Studio gives warnings when using
+// fopen. But fopen_s is not going to work well
+// with most compilers, and fopen_s uses different
+// syntax than fopen. (i.e., a macro won't work)
+// So, we'lll use this:
+#define _CRT_SECURE_NO_DEPRECATE
+#endif
+
+#include <iostream>
+#include <cmath>
+#include <cctype>
+#include <cstring>
+
+#ifndef EasyBMP
+#define EasyBMP
+
+#ifdef __BCPLUSPLUS__
+// The Borland compiler must use this because something
+// is wrong with their cstdio file.
+#include <stdio.h>
+#else
+#include <cstdio>
+#endif
+
+#ifdef __GNUC__
+// If g++ specific code is ever required, this is
+// where it goes.
+#endif
+
+#ifdef __INTEL_COMPILER
+// If Intel specific code is ever required, this is
+// where it goes.
+#endif
+
+#ifndef _DefaultXPelsPerMeter_
+#define _DefaultXPelsPerMeter_
+#define DefaultXPelsPerMeter 3780
+// set to a default of 96 dpi
+#endif
+
+#ifndef _DefaultYPelsPerMeter_
+#define _DefaultYPelsPerMeter_
+#define DefaultYPelsPerMeter 3780
+// set to a default of 96 dpi
+#endif
+
+#include "EasyBMP_DataStructures.h"
+#include "EasyBMP_BMP.h"
+#include "EasyBMP_VariousBMPutilities.h"
+
+#ifndef _EasyBMP_Version_
+#define _EasyBMP_Version_ 1.06
+#define _EasyBMP_Version_Integer_ 106
+#define _EasyBMP_Version_String_ "1.06"
+#endif
+
+#ifndef _EasyBMPwarnings_
+#define _EasyBMPwarnings_
+#endif
+
+void SetEasyBMPwarningsOff( void );
+void SetEasyBMPwarningsOn( void );
+bool GetEasyBMPwarningState( void );
+
+#endif
diff --git a/benchmarks/CUDA/RAY/EasyBMP_BMP.h b/benchmarks/CUDA/RAY/EasyBMP_BMP.h
new file mode 100644
index 0000000..819a976
--- /dev/null
+++ b/benchmarks/CUDA/RAY/EasyBMP_BMP.h
@@ -0,0 +1,86 @@
+/*************************************************
+* *
+* EasyBMP Cross-Platform Windows Bitmap Library *
+* *
+* Author: Paul Macklin *
+* email: [email protected] *
+* support: http://easybmp.sourceforge.net *
+* *
+* file: EasyBMP_VariousBMPutilities.h *
+* date added: 05-02-2005 *
+* date modified: 12-01-2006 *
+* version: 1.06 *
+* *
+* License: BSD (revised/modified) *
+* Copyright: 2005-6 by the EasyBMP Project *
+* *
+* description: Defines BMP class *
+* *
+*************************************************/
+
+#ifndef _EasyBMP_BMP_h_
+#define _EasyBMP_BMP_h_
+
+bool SafeFread( char* buffer, int size, int number, FILE* fp );
+bool EasyBMPcheckDataSize( void );
+
+class BMP
+{private:
+
+ int BitDepth;
+ int Width;
+ int Height;
+ RGBApixel** Pixels;
+ RGBApixel* Colors;
+ int XPelsPerMeter;
+ int YPelsPerMeter;
+
+ ebmpBYTE* MetaData1;
+ int SizeOfMetaData1;
+ ebmpBYTE* MetaData2;
+ int SizeOfMetaData2;
+
+ bool Read32bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+ bool Read24bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+ bool Read8bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+ bool Read4bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+ bool Read1bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+
+ bool Write32bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+ bool Write24bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+ bool Write8bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+ bool Write4bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+ bool Write1bitRow( ebmpBYTE* Buffer, int BufferSize, int Row );
+
+ ebmpBYTE FindClosestColor( RGBApixel& input );
+
+ public:
+
+ int TellBitDepth( void );
+ int TellWidth( void );
+ int TellHeight( void );
+ int TellNumberOfColors( void );
+ void SetDPI( int HorizontalDPI, int VerticalDPI );
+ int TellVerticalDPI( void );
+ int TellHorizontalDPI( void );
+
+ BMP();
+ BMP( BMP& Input );
+ ~BMP();
+ RGBApixel* operator()(int i,int j);
+
+ RGBApixel GetPixel( int i, int j ) const;
+ bool SetPixel( int i, int j, RGBApixel NewPixel );
+
+ bool CreateStandardColorTable( void );
+
+ bool SetSize( int NewWidth, int NewHeight );
+ bool SetBitDepth( int NewDepth );
+ bool WriteToFile( const char* FileName );
+ bool ReadFromFile( const char* FileName );
+
+ RGBApixel GetColor( int ColorNumber );
+ bool SetColor( int ColorNumber, RGBApixel NewColor );
+};
+
+#endif
diff --git a/benchmarks/CUDA/RAY/EasyBMP_DataStructures.h b/benchmarks/CUDA/RAY/EasyBMP_DataStructures.h
new file mode 100644
index 0000000..82b6179
--- /dev/null
+++ b/benchmarks/CUDA/RAY/EasyBMP_DataStructures.h
@@ -0,0 +1,104 @@
+/*************************************************
+* *
+* EasyBMP Cross-Platform Windows Bitmap Library *
+* *
+* Author: Paul Macklin *
+* email: [email protected] *
+* support: http://easybmp.sourceforge.net *
+* *
+* file: EasyBMP_DataStructures.h *
+* date added: 05-02-2005 *
+* date modified: 12-01-2006 *
+* version: 1.06 *
+* *
+* License: BSD (revised/modified) *
+* Copyright: 2005-6 by the EasyBMP Project *
+* *
+* description: Defines basic data structures for *
+* the BMP class *
+* *
+*************************************************/
+
+#ifndef _EasyBMP_Custom_Math_Functions_
+#define _EasyBMP_Custom_Math_Functions_
+inline double Square( double number )
+{ return number*number; }
+
+inline int IntSquare( int number )
+{ return number*number; }
+#endif
+
+int IntPow( int base, int exponent );
+
+#ifndef _EasyBMP_Defined_WINGDI
+#define _EasyBMP_Defined_WINGDI
+ typedef unsigned char ebmpBYTE;
+ typedef unsigned short ebmpWORD;
+ typedef unsigned int ebmpDWORD;
+#endif
+
+#ifndef _EasyBMP_DataStructures_h_
+#define _EasyBMP_DataStructures_h_
+
+inline bool IsBigEndian()
+{
+ short word = 0x0001;
+ if((*(char *)& word) != 0x01 )
+ { return true; }
+ return false;
+}
+
+inline ebmpWORD FlipWORD( ebmpWORD in )
+{ return ( (in >> 8) | (in << 8) ); }
+
+inline ebmpDWORD FlipDWORD( ebmpDWORD in )
+{
+ return ( ((in&0xFF000000)>>24) | ((in&0x000000FF)<<24) |
+ ((in&0x00FF0000)>>8 ) | ((in&0x0000FF00)<<8 ) );
+}
+
+// it's easier to use a struct than a class
+// because we can read/write all four of the bytes
+// at once (as we can count on them being continuous
+// in memory
+
+typedef struct RGBApixel {
+ ebmpBYTE Blue;
+ ebmpBYTE Green;
+ ebmpBYTE Red;
+ ebmpBYTE Alpha;
+} RGBApixel;
+
+class BMFH{
+public:
+ ebmpWORD bfType;
+ ebmpDWORD bfSize;
+ ebmpWORD bfReserved1;
+ ebmpWORD bfReserved2;
+ ebmpDWORD bfOffBits;
+
+ BMFH();
+ void display( void );
+ void SwitchEndianess( void );
+};
+
+class BMIH{
+public:
+ ebmpDWORD biSize;
+ ebmpDWORD biWidth;
+ ebmpDWORD biHeight;
+ ebmpWORD biPlanes;
+ ebmpWORD biBitCount;
+ ebmpDWORD biCompression;
+ ebmpDWORD biSizeImage;
+ ebmpDWORD biXPelsPerMeter;
+ ebmpDWORD biYPelsPerMeter;
+ ebmpDWORD biClrUsed;
+ ebmpDWORD biClrImportant;
+
+ BMIH();
+ void display( void );
+ void SwitchEndianess( void );
+};
+
+#endif
diff --git a/benchmarks/CUDA/RAY/EasyBMP_VariousBMPutilities.h b/benchmarks/CUDA/RAY/EasyBMP_VariousBMPutilities.h
new file mode 100644
index 0000000..349dda6
--- /dev/null
+++ b/benchmarks/CUDA/RAY/EasyBMP_VariousBMPutilities.h
@@ -0,0 +1,43 @@
+/*************************************************
+* *
+* EasyBMP Cross-Platform Windows Bitmap Library *
+* *
+* Author: Paul Macklin *
+* email: [email protected] *
+* support: http://easybmp.sourceforge.net *
+* *
+* file: EasyBMP_VariousBMPutilities.h *
+* date added: 05-02-2005 *
+* date modified: 12-01-2006 *
+* version: 1.06 *
+* *
+* License: BSD (revised/modified) *
+* Copyright: 2005-6 by the EasyBMP Project *
+* *
+* description: Various utilities. *
+* *
+*************************************************/
+
+#ifndef _EasyBMP_VariousBMPutilities_h_
+#define _EasyBMP_VariousBMPutilities_h_
+
+BMFH GetBMFH( const char* szFileNameIn );
+BMIH GetBMIH( const char* szFileNameIn );
+void DisplayBitmapInfo( const char* szFileNameIn );
+int GetBitmapColorDepth( const char* szFileNameIn );
+void PixelToPixelCopy( BMP& From, int FromX, int FromY,
+ BMP& To, int ToX, int ToY);
+void PixelToPixelCopyTransparent( BMP& From, int FromX, int FromY,
+ BMP& To, int ToX, int ToY,
+ RGBApixel& Transparent );
+void RangedPixelToPixelCopy( BMP& From, int FromL , int FromR, int FromB, int FromT,
+ BMP& To, int ToX, int ToY );
+void RangedPixelToPixelCopyTransparent(
+ BMP& From, int FromL , int FromR, int FromB, int FromT,
+ BMP& To, int ToX, int ToY ,
+ RGBApixel& Transparent );
+bool CreateGrayscaleColorTable( BMP& InputImage );
+
+bool Rescale( BMP& InputImage , char mode, int NewDimension );
+
+#endif
diff --git a/benchmarks/CUDA/RAY/Makefile b/benchmarks/CUDA/RAY/Makefile
new file mode 100644
index 0000000..b83a221
--- /dev/null
+++ b/benchmarks/CUDA/RAY/Makefile
@@ -0,0 +1,50 @@
+################################################################################
+#
+# Copyright 1993-2006 NVIDIA Corporation. All rights reserved.
+#
+# NOTICE TO USER:
+#
+# This source code is subject to NVIDIA ownership rights under U.S. and
+# international Copyright laws.
+#
+# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
+# CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
+# IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH
+# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
+# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
+# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,
+# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS
+# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE
+# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE
+# OR PERFORMANCE OF THIS SOURCE CODE.
+#
+# U.S. Government End Users. This source code is a "commercial item" as
+# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of
+# "commercial computer software" and "commercial computer software
+# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995)
+# and is provided to the U.S. Government only as a commercial end item.
+# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
+# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
+# source code with only those rights set forth herein.
+#
+################################################################################
+#
+# Build script for project
+#
+################################################################################
+
+# Add source files here
+EXECUTABLE := ray_2
+# Cuda source files (compiled with cudacc)
+CUFILES += rayTracing.cu
+CCFILES := EasyBMP.cpp makebmp.cpp
+
+GPGPUSIM_ROOT := ../../..
+
+################################################################################
+#USEGLLIB := 1
+
+# Rules and targets
+#LIB :=
+
+include ../../../common/common.mk
diff --git a/benchmarks/CUDA/RAY/README.GPGPU-Sim b/benchmarks/CUDA/RAY/README.GPGPU-Sim
new file mode 100644
index 0000000..06462ed
--- /dev/null
+++ b/benchmarks/CUDA/RAY/README.GPGPU-Sim
@@ -0,0 +1,2 @@
+make
+./gpgpu_ptx_sim__ray_2 256 256
diff --git a/benchmarks/CUDA/RAY/makebmp.cpp b/benchmarks/CUDA/RAY/makebmp.cpp
new file mode 100644
index 0000000..1b006ec
--- /dev/null
+++ b/benchmarks/CUDA/RAY/makebmp.cpp
@@ -0,0 +1,28 @@
+#include "EasyBMP.h"
+
+BMP out_bmp;
+
+void initialize_bmp(unsigned width, unsigned height, unsigned depth)
+{
+ SetEasyBMPwarningsOff();
+ out_bmp.SetSize(width, height);
+ out_bmp.SetBitDepth(depth);
+}
+
+void create_bmp(unsigned *data)
+{
+ unsigned height = out_bmp.TellHeight();
+ unsigned width = out_bmp.TellWidth();
+ for (unsigned y=0; y< height; y++){
+ for (unsigned x=0; x< width; x++) {
+ //printf("%8x ", c_output[x+y*width]);
+ out_bmp(x,(height-y-1))->Red = 0x000000FF & data[x+y*width];
+ out_bmp(x,(height-y-1))->Green = (0x0000FF00 & data[x+y*width]) >> 8;
+ out_bmp(x,(height-y-1))->Blue = (0x00FF0000 & data[x+y*width]) >> 16;
+ out_bmp(x,(height-y-1))->Alpha = (0xFF000000 & data[x+y*width]) >> 24;
+ }
+ }
+ out_bmp.WriteToFile("output.bmp");
+}
+
+
diff --git a/benchmarks/CUDA/RAY/makebmp.h b/benchmarks/CUDA/RAY/makebmp.h
new file mode 100644
index 0000000..1aabfcc
--- /dev/null
+++ b/benchmarks/CUDA/RAY/makebmp.h
@@ -0,0 +1,9 @@
+#ifndef _MAKEBMP_H_
+#define _MAKEBMP_H_
+
+
+void initialize_bmp(unsigned width, unsigned height, unsigned depth);
+
+void create_bmp(unsigned *data);
+
+#endif // _MAKEBMP_H_
diff --git a/benchmarks/CUDA/RAY/rayTracing.cu b/benchmarks/CUDA/RAY/rayTracing.cu
new file mode 100644
index 0000000..80145dd
--- /dev/null
+++ b/benchmarks/CUDA/RAY/rayTracing.cu
@@ -0,0 +1,441 @@
+/*
+ * Copyright 2008 BOROUJERDI Maxime. Tous droits reserves.
+ */
+
+#include <cstdlib>
+#include <cstdio>
+#include <cstring>
+#include <cmath>
+#include "makebmp.h"
+
+/*#include <GL/glew.h>
+#include <GL/glut.h>
+
+#include <cuda_gl_interop.h>*/
+#include <cutil.h>
+
+typedef unsigned int uint;
+typedef unsigned char uchar;
+
+#define numObj 4
+
+#define PI 3.141592654f
+#define Angle(a) ((a*PI)/180.0)
+
+//#define DEVICE_EMU
+//#define DEBUG_RT_CUDA
+#define FIXED_CONST_PARSE
+#ifdef DEBUG_RT_CUDA
+#define DEBUG_NUM 8
+float4 *d_debug_float4;
+uint *d_debug_uint;
+float4 *h_debug_float4;
+uint *h_debug_uint;
+#endif
+int g_verbose;
+
+#include <rayTracing_kernel.cu>
+
+unsigned width = 64; //640; //512; //16; //32; //512;
+unsigned height = 64; //480; //512; //16;//512;
+dim3 blockSize(16,8);
+dim3 gridSize(width/blockSize.x, height/blockSize.y);
+
+float3 viewRotation;
+float3 viewTranslation = make_float3(0.0, 0.0, -4.0f);
+float invViewMatrix[12];
+
+//static int fpsCount = 0; // FPS count for averaging
+//static int fpsLimit = 1; // FPS limit for sampling
+unsigned int timer;
+
+
+//GLuint pbo = 0; // Pixel buffer d'OpenGL
+
+
+void initPixelBuffer();
+
+class Observateur
+{
+ private:
+ matrice3x4 M; // U, V, W
+ float df; // distance focale
+
+ public:
+ Observateur( );
+ Observateur(const float3 &, const float3 &, const float3 &, double );
+
+ inline const matrice3x4 & getMatrice( ) const { return M; }
+ inline float getDistance( ) const { return df; }
+};
+
+Observateur::Observateur()
+{
+ M.m[0] = make_float4(0.0f,0.0f,1.0f,0.0f);
+ M.m[1] = make_float4(0.0f,1.0f,0.0f,0.0f);
+ M.m[2] = make_float4(1.0f,0.0f,0.0f,0.0f);
+ df = 1.0 / tan(Angle(65)/2.0);
+}
+
+Observateur::Observateur(const float3 & p, const float3 & u, const float3 & v, double a )
+{
+ float3 VP, U, V, W;
+ VP = normalize(v);
+ U = normalize(u);
+ V = normalize(VP - dot(U,VP)*U);
+ W = normalize(cross(U,V));
+ M.m[0] = make_float4(U.x,U.y,U.z,p.x);
+ M.m[1] = make_float4(V.x,V.y,V.z,p.y);
+ M.m[2] = make_float4(W.x,W.y,W.z,p.z);
+ df = 1.0 / tan(Angle(a)/2.0);
+}
+
+float anim = 0.0f, pas = 0.015f;
+Observateur obs = Observateur(make_float3(0.0f,0.5f,2.0f),normalize(make_float3(0.0f,0.0f,0.0f)-make_float3(0.0f,0.5f,2.0f)),make_float3(0.0f,1.0f,0.0f),65.0f);;
+
+uint * values = NULL, * d_output, * d_temp, NUM;
+uint * c_output;
+
+Node node[numObj], * d_node;
+
+Sphere s, s1, s2;
+float phi;
+
+uint * nObj;
+float * prof;
+Rayon * ray;
+float3 * A, *u;
+int t = 1;
+
+
+void initObjet()
+{
+ srand(47);
+ node->s.r = 1.0f;
+ node[0].s.C = make_float3(0.0f,-1.5f,-0.0f); node[0].s.r = 0.5f;
+ node[1].s.C = make_float3(-1.0f,0.0f,-1.0f); node[1].s.r = 0.5f;
+ node[2].s.C = make_float3(1.0f,-0.f,-1.0f); node[2].s.r = 0.5f;
+ node[3].s.C = make_float3(0.0f,-0.f,-2.0f); node[3].s.r = 0.75f;
+ for( int i(4); i < numObj; i++ ) {
+ float r,v,b;
+ float tmp1(5.0f*((r=(float(rand()%255)/255.0f)))-2.5f);
+ float tmp2(5.0f*((v=(float(rand()%255)/255.0f)))-2.5f);
+ float tmp3(-5.0f*((b=(float(rand()%255)/255.0f))));
+ float tmp4((rand()%100)/100.0f);
+ node[i].s.C = make_float3(tmp1,tmp2,tmp3); node[i].s.r = tmp4;
+ node[i].s.R = r; node[i].s.V = v; node[i].s.B = b; node[i].s.A = 1.0f;
+ node[i].fg = 0; node[i].fd = 0;
+ }
+ node[0].s.R = 0.0f; node[0].s.V = 1.0f; node[0].s.B = 1.0f; node[0].s.A = 1.0f;
+ node[1].s.R = 1.0f; node[1].s.V = 0.0f; node[1].s.B = 0.0f; node[1].s.A = 1.0f;
+ node[2].s.R = 0.0f; node[2].s.V = 0.0f; node[2].s.B = 1.0f; node[2].s.A = 1.0f;
+ node[3].s.R = 0.0f; node[3].s.V = 1.0f; node[3].s.B = 0.0f; node[3].s.A = 1.0f;
+ //createNode(&node[0], &node[1], &node[2], 1.0f);
+ node[0].fg = 1; node[0].fd = 2;
+ node[1].fg = 0; node[1].fd = 0;
+ node[2].fg = 0; node[2].fd = 0;
+ node[3].fg = 0; node[3].fd = 0;
+
+ #ifdef DEBUG_RT_CUDA
+ h_debug_float4 = (float4*) calloc(DEBUG_NUM, sizeof(float4));
+ h_debug_uint = (uint*) calloc(DEBUG_NUM, sizeof(uint));
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&d_debug_float4, DEBUG_NUM*sizeof(float4)));
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&d_debug_uint, DEBUG_NUM*sizeof(uint)));
+ CUDA_SAFE_CALL( cudaMemcpy( d_debug_float4, h_debug_float4, DEBUG_NUM*sizeof(float4), cudaMemcpyHostToDevice) );
+ CUDA_SAFE_CALL( cudaMemcpy( d_debug_uint, h_debug_uint, DEBUG_NUM*sizeof(uint), cudaMemcpyHostToDevice) );
+ #endif
+ c_output = (uint*) calloc(width*height, sizeof(uint));
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&d_output, width*height*sizeof(uint)));
+
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&d_node, numObj*sizeof(Node) ));
+ CUDA_SAFE_CALL( cudaMemcpy( d_node, node, numObj*sizeof(Node), cudaMemcpyHostToDevice) );
+ CUDA_SAFE_CALL( cudaMemcpyToSymbol(cnode, node, numObj*sizeof(Node)) );
+ CUDA_SAFE_CALL( cudaMemcpyToSymbol(MView, (void*)&obs, 3*sizeof(float4)) );
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&d_temp, width * height*sizeof(uint)));
+ CUDA_SAFE_CALL( cudaMemset(d_temp, 0, width * height*sizeof(uint)) );
+
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&nObj, width * height*sizeof(uint)));
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&prof, width * height*sizeof(float)));
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&ray, width * height*sizeof(Rayon)));
+
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&A, width * height*sizeof(float3)));
+ CUDA_SAFE_CALL( cudaMalloc( (void**)&u, width * height*sizeof(float3)));
+}
+
+#define PRINT_PIXELS
+
+// Rendu de l'image avec CUDA
+void render()
+{
+ // map PBO to get CUDA device pointer <GY: replace with memcpy?>
+ //CUDA_SAFE_CALL(cudaGLMapBufferObject((void**)&d_output, pbo));
+ //CUDA_SAFE_CALL( cudaMemcpy( d_output, c_output, width*height*sizeof(uint), cudaMemcpyHostToDevice) );
+ // call CUDA kernel, writing results to PBO
+ CUT_SAFE_CALL(cutStartTimer(timer));
+ #ifdef DEBUG_RT_CUDA
+ render<<<gridSize, blockSize>>>(d_debug_float4, d_debug_uint, d_output, d_node, width, height, anim, obs.getDistance());
+ #else
+ render<<<gridSize, blockSize>>>(d_output, d_node, width, height, anim, obs.getDistance());
+ #endif
+ CUDA_SAFE_CALL( cudaThreadSynchronize() );
+ CUT_SAFE_CALL(cutStopTimer(timer));
+
+ #ifdef DEBUG_RT_CUDA
+ CUDA_SAFE_CALL( cudaMemcpy( h_debug_float4, d_debug_float4, DEBUG_NUM*sizeof(float4), cudaMemcpyDeviceToHost) );
+ CUDA_SAFE_CALL( cudaMemcpy( h_debug_uint, d_debug_uint, DEBUG_NUM*sizeof(uint), cudaMemcpyDeviceToHost) );
+
+ printf("debug_float4\n");
+ for (int i=0; i< DEBUG_NUM; i++) {
+ printf("%e %e %e %e\n", h_debug_float4[i].x, h_debug_float4[i].y, h_debug_float4[i].z, h_debug_float4[i].w);
+ }
+ printf("debug_uint\n");
+ for (int i=0; i< DEBUG_NUM; i++) {
+ printf("0x%x\n", h_debug_uint[i]);
+ }
+ #endif
+
+ CUDA_SAFE_CALL( cudaMemcpy( c_output, d_output, width*height*sizeof(uint), cudaMemcpyDeviceToHost) );
+ unsigned long long int checksum = 0;
+ for (int y=(height-1); y >= 0; y--){
+ if (g_verbose) printf("\n");
+ for (int x=0; x< width; x++) {
+ if (g_verbose) printf("%010u ", (unsigned) c_output[x+y*width]);
+ checksum += c_output[x+y*width];
+ }
+ }
+ printf("\n");
+ printf("checksum=%llx\n", checksum);
+ CUT_CHECK_ERROR("Erreur kernel");
+
+ //CUDA_SAFE_CALL(cudaGLUnmapBufferObject(pbo)); //<GY: replace with memcpy?>
+
+}
+
+// Affichage du resultat avec OpenGL
+void display()
+{
+
+ // Affichage du resultat
+ //glClear(GL_COLOR_BUFFER_BIT);
+
+ //CUT_SAFE_CALL(cutStartTimer(timer));
+ render();
+ //CUT_SAFE_CALL(cutStopTimer(timer));
+ printf("Kernel Time: %f \n", cutGetTimerValue(timer));
+ /*fpsCount++;
+ if (fpsCount == fpsLimit) {
+ char fps[256];
+ float ifps = 1.f / (cutGetAverageTimerValue(timer) / 1000.f);
+ sprintf(fps, "Cuda Ray Tracing: %.1f fps", ifps);
+ glutSetWindowTitle(fps);
+ fpsCount = 0;
+ fpsLimit = (int)max(ifps, 1.f);
+ CUT_SAFE_CALL(cutResetTimer(timer));
+ }*/
+
+ if( anim >= 1.0f ) pas = -0.015f;
+ else if( anim <= -1.0f ) pas = 0.015f;
+ anim += pas;
+
+ // Dessin de l'image de PBO
+ /*glDisable(GL_DEPTH_TEST);
+ glRasterPos2i(0, 0);
+ glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, pbo);
+ glDrawPixels(width, height, GL_RGBA, GL_UNSIGNED_BYTE, 0);
+ glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
+
+ glutSwapBuffers();
+ glutReportErrors();*/
+ t--;
+ if (!t) {
+ return;
+ }
+
+}
+
+/*void idle()
+{
+ glutPostRedisplay();
+}
+
+void keyboard(unsigned char , int , int )
+{
+ //glutPostRedisplay();
+}*/
+
+int ox, oy;
+int buttonState = 0;
+
+/*void mouse(int , int , int , int )
+{
+ if (state == GLUT_DOWN)
+ buttonState |= 1<<button;
+ else if (state == GLUT_UP)
+ buttonState = 0;
+
+ ox = x; oy = y;
+ glutPostRedisplay();
+}
+
+void motion(int , int )
+{
+ float dx, dy;
+ dx = x - ox;
+ dy = y - oy;
+
+ if (buttonState == 3) {
+ // left+middle = zoom
+ viewTranslation.z += dy / 100.0;
+ }
+ else if (buttonState & 2) {
+ // middle = translate
+ viewTranslation.x += dx / 100.0;
+ viewTranslation.y -= dy / 100.0;
+ }
+ else if (buttonState & 1) {
+ // left = rotate
+ viewRotation.x += dy / 5.0;
+ viewRotation.y += dx / 5.0;
+ }
+
+ ox = x; oy = y;
+ glutPostRedisplay();
+}
+
+void reshape(int x, int y)
+{
+ width = x; height = y;
+ initPixelBuffer();
+
+ glViewport(0, 0, x, y);
+ //glViewport(-x/2, -y/2, x/2, y/2);
+ glMatrixMode(GL_MODELVIEW);
+ glLoadIdentity();
+
+ glMatrixMode(GL_PROJECTION);
+ glLoadIdentity();
+ glOrtho(0.0, 1.0, 0.0, 1.0, 0.0, 1.0);
+}
+
+void cleanup()
+{
+ CUDA_SAFE_CALL(cudaGLUnregisterBufferObject(pbo));
+ glDeleteBuffersARB(1, &pbo);
+ CUT_SAFE_CALL(cutDeleteTimer(timer));
+}*/
+
+int iDivUp(int a, int b)
+{
+ return (a % b != 0) ? (a / b + 1) : (a / b);
+}
+
+void initPixelBuffer()
+{
+ /*if (pbo) {
+ // delete old buffer
+ CUDA_SAFE_CALL(cudaGLUnregisterBufferObject(pbo));
+ glDeleteBuffersARB(1, &pbo);
+ }*/
+
+ NUM = width * height;
+ phi = 2.0f/(float)min(width,height);
+
+ // create pixel buffer object for display
+ /* glGenBuffersARB(1, &pbo);
+ glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, pbo);
+ glBufferDataARB(GL_PIXEL_UNPACK_BUFFER_ARB, width*height*sizeof(GLubyte)*4, 0, GL_STREAM_DRAW_ARB);
+ glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
+
+ CUDA_SAFE_CALL(cudaGLRegisterBufferObject(pbo));*/
+
+ // calculate new grid size
+ gridSize = dim3(iDivUp(width, blockSize.x), iDivUp(height, blockSize.y));
+}
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Programme principal
+////////////////////////////////////////////////////////////////////////////////
+
+
+
+int main( int argc, char** argv)
+{
+ // initialise card and timer
+ int deviceCount;
+ CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceCount(&deviceCount));
+ if (deviceCount == 0) {
+ fprintf(stderr, "There is no device.\n");
+ exit(EXIT_FAILURE);
+ }
+ int dev;
+ for (dev = 0; dev < deviceCount; ++dev) {
+ cudaDeviceProp deviceProp;
+ CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceProperties(&deviceProp, dev));
+ if (deviceProp.major >= 1)
+ break;
+ }
+ if (dev == deviceCount) {
+ fprintf(stderr, "There is no device supporting CUDA.\n");
+ exit(EXIT_FAILURE);
+ }
+ else
+ CUDA_SAFE_CALL(cudaSetDevice(dev));
+ int i, commandline_error;
+ commandline_error = 0;
+ g_verbose = 0;
+ if (argc >= 3) {
+ width = atoi(argv[1]);
+ height = atoi(argv[2]);
+ for (i=3; i < argc;i++) {
+ if (argv[i][0] == '-') {
+ switch (argv[i][1]) {
+ case 'v': g_verbose = 1;
+ break;
+ default: commandline_error=1;
+ }
+ }
+ else commandline_error=1;
+ }
+ } else commandline_error=1;
+
+ if (commandline_error || !width || !height) {
+ printf("Usage: ./rayTracing <WIDTH> <HEIGHT> [-v]\n");
+ printf("where WIDTH and HEIGHT are the screen dimensions and -v is used to display an abstract representation of the output.\n");
+ return 1;
+ }
+ CUT_SAFE_CALL(cutCreateTimer(&timer));
+ CUT_SAFE_CALL(cutResetTimer(timer));
+
+ initialize_bmp(width,height,32);
+
+ // initialise les functions callback de GLUT
+ /*glutInit(&argc, argv);
+ glutInitDisplayMode(GLUT_RGB | GLUT_DOUBLE);
+ glutInitWindowSize(width, height);
+ glutCreateWindow("CUDA Ray Tracing");
+ glutDisplayFunc(display);
+ glutKeyboardFunc(keyboard);
+ glutMouseFunc(mouse);
+ glutMotionFunc(motion);
+ glutReshapeFunc(reshape);
+ glutIdleFunc(idle);
+
+ glewInit();
+ if (!glewIsSupported("GL_VERSION_2_0 GL_ARB_pixel_buffer_object")) {
+ fprintf(stderr, "Les extensions minimales d'OpenGL sont absentes.");
+ exit(-1);
+ }
+ initPixelBuffer();
+ initObjet();
+
+ atexit(cleanup);
+
+ glutMainLoop();*/
+ initObjet();
+ initPixelBuffer();
+ display();
+ create_bmp(c_output);
+ CUT_SAFE_CALL(cutDeleteTimer(timer));
+ return 0;
+}
diff --git a/benchmarks/CUDA/RAY/rayTracing_kernel.cu b/benchmarks/CUDA/RAY/rayTracing_kernel.cu
new file mode 100644
index 0000000..c524a9c
--- /dev/null
+++ b/benchmarks/CUDA/RAY/rayTracing_kernel.cu
@@ -0,0 +1,658 @@
+/*
+ * Copyright 2008 BOROUJERDI Maxime. Tous droits reserves.
+ */
+
+//#define FIXED_CONST_PARSE
+#ifndef __RAYTRACING_KERNEL_H__
+#define __RAYTRACING_KERNEL_H__
+
+#include "cutil_math.h"
+
+typedef struct
+{
+ float4 m[3];
+} matrice3x4;
+
+typedef struct {
+ float4 m[4];
+} matrice4x4;
+
+typedef struct{
+ float3 A; // origine
+ float3 u; // direction
+} Rayon;
+
+typedef struct Sphere{
+ float3 C; // centre
+ float r; // rayon
+ float R,V,B,A;
+ /*Sphere() : C(make_float3(0.0f,0.0f,0.0f)), r(0.5f), rvba(make_float4(1.0f,0.0f,0.0f,1.0f)) { }
+ Sphere(const float3 _C, float _r, const float4 _rvba) : C(_C), r(_r), rvba(_rvba) { }
+ Sphere(const float3 _C, float _r) : C(_C), r(_r), rvba(make_float4(1.0f,0.0f,0.0f,1.0f)) { }*/
+} Sphere;
+
+typedef struct Node {
+ Sphere s;
+ uint fg, fd;
+} Node;
+
+/*__host__ __device__ void createNode(Node * n, Node * fg, Node * fd, const Sphere & s)
+{
+ n->fg = fg;
+ n->fd = fd;
+ n->C = s.C;
+ n->r = s.r;
+}
+
+__host__ __device__ Node * filsGauche(Node * n) { return n->fg; }
+
+__host__ __device__ Node * filsDroite(Node * n) { return n->fd; }*/
+
+//__host__ __device__ Sphere sphere(Node * n) { return n->s; }
+
+__constant__ matrice3x4 MView; // matrice inverse de la matrice de vue
+
+__constant__ Node cnode[numObj];
+
+template <class T>
+__device__ void swap(T & v1, T & v2)
+{
+ T tmp(v1);
+ v1 = v2;
+ v2 = tmp;
+}
+
+__device__ float intersectionSphere(Rayon R, float3 C, float r)
+{
+ float3 L(C-R.A);
+ float d(dot(L,R.u)), l2(dot(L,L)), r2(r*r), m2, q, res;
+
+ if( d < 0.0f && l2 > r2 ) {
+ res = 0.0f;
+ }
+ else
+ {
+ m2 = l2 - d*d;
+ if( m2 > r2 ) {
+ res = 0.0f;
+ }
+ else
+ {
+ q = sqrt(r2-m2);
+ if( l2 > r2 ) res = d - q;
+ else res = d + q;
+ }
+ }
+
+ return res;
+}
+
+__device__ float intersectionPlan( Rayon R, float3 C, float3 N2 )
+{
+ float res;
+ float3 N = normalize(make_float3(0.0f,1.0f,0.0f));
+ float m(dot(N,R.u)), d, t;
+ float3 L;
+
+ if( fabs(m) < 0.0001f ) {
+ res = 0.0f;
+ }
+ else {
+ L = R.A - C;
+ d = dot(N,L);
+ t = -d/m;
+ if( t > 0 ) {
+ res = t;
+ }
+ else {
+ res = 0.0f;
+ }
+ }
+
+ return res;
+}
+
+__device__ float3 getNormale(float3 P, float3 C)
+{
+ return normalize(P-C);
+}
+
+__device__ float3 getNormaleP(float3 P)
+{
+ return normalize(make_float3(0.0f,1.0f,0.0f));
+}
+
+// multiplication d'un vecteur par une matrice (sans translation)
+__device__ float3 mul(matrice3x4 M, float3 v)
+{
+ float3 r;
+ r.x = dot(v, make_float3(M.m[0]));
+ r.y = dot(v, make_float3(M.m[1]));
+ r.z = dot(v, make_float3(M.m[2]));
+ return r;
+}
+
+// multiplication d'un vecteur par une matrice avec translation
+__device__ float4 mul(matrice3x4 M, float4 v)
+{
+ float4 r;
+ r.x = dot(v, M.m[0]);
+ r.y = dot(v, M.m[1]);
+ r.z = dot(v, M.m[2]);
+ r.w = 1.0f;
+ return r;
+}
+
+__device__ uint rgbaFloatToInt(float4 rgba)
+{
+ #ifdef DEVICE_EMU
+ printf("%d: rgba = %f %f %f %f\n", threadIdx.x, rgba.x, rgba.y, rgba.z, rgba.w);
+ #endif
+ rgba.x = __saturatef(rgba.x); // clamp entre [0.0, 1.0]
+ rgba.y = __saturatef(rgba.y);
+ rgba.z = __saturatef(rgba.z);
+ rgba.w = __saturatef(rgba.w);
+#ifdef DEVICE_EMU
+ printf("%d: rgba = %x %x %x %x\n", threadIdx.x, uint(rgba.x*255), uint(rgba.y*255), uint(rgba.z*255), uint(rgba.w*255));
+#endif
+ return (uint(rgba.w*255)<<24)
+ | (uint(rgba.z*255)<<16)
+ | (uint(rgba.y*255)<<8 )
+ | (uint(rgba.x*255) );
+}
+
+/*__device__ void myswap(Sphere &x, Sphere &y)
+{
+Sphere t = x;
+x = y;
+y = t;
+}*/
+/*__global__ void d_render(uint * d_output, uint imageW, uint imageH, float pas, float df, float tPixel)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint id = x + y * imageW;
+
+ if( x < imageW && y < imageH )
+ {
+ //float tPixel = 2.0f/(float)min(imageW,imageH);
+ matrice3x4 M(MView);
+ Rayon R;
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+ Sphere s(cnode[1].s), s2(cnode[2].s), st(cnode[2].s);
+ float t, t2, tt;
+ s.C.x += pas, s2.C.x += pas;
+ t = intersectionSphere(R,s.C,s.r);
+ t2 = intersectionSphere(R,s2.C,s2.r);
+ if( !t ) {
+ //myswap(s,s2);
+ //swap(t,t2);
+ tt = t;
+ t = t2;
+ t2 = tt;
+ st = s;
+ s = s2;
+ s2 = st;
+ }
+ else if( t2 && t2 < t ) {
+ //myswap(s,s2);
+ //swap(t,t2);
+ tt = t;
+ t = t2;
+ t2 = tt;
+ st = s;
+ s = s2;
+ s2 = st;
+ }
+ float4 f = make_float4(0,1,0,1)*(dot(getNormale(R.A+R.u*t,s.C),(-1.0f)*R.u));
+ uint n = rgbaFloatToInt(f);
+ //printf("%f\n",d_node[0].s.r);
+ if( t > 0.0f )
+ d_output[id] = n;
+ //else d_output[id] = 0;
+ }
+ __syncthreads();
+}
+*/
+/*__global__ void rayCast (uint * d_output, uint * d_temp, uint imageW, uint imageH, float pas, float df)
+//(uint * result, uint * temp, uint imageW, uint imageH, float pas, float df)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y);
+ uint id = x + y * gridDim.x;
+ //float tmp= float(imageW)/float(gridDim.x);
+ float t;
+
+ //if( x < gridDim.x && y < gridDim.y )
+ if( d_temp[id] == 0 )
+ {
+ float tPixel = 2.0f/float(imageW);
+ matrice3x4 M(MView);
+ Rayon R;
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+ Sphere s(cnode[1].s);
+ s.C.x += pas;
+ t = intersectionSphere(R,s.C,s.r/(imageW/gridDim.x));
+
+ if( t > 0.0f ) {
+ //float4 f = make_float4(0,1,0,1)*(dot(getNormale(R.A+R.u*t,s.C),(-1.0f)*R.u));
+ d_output[id] = rgbaFloatToInt(make_float4(0,1,0,1));
+ //printf("%d %d\n",int(x*tmp),int((y*tmp)/2));
+ }
+ else {
+// float tmp= float(imageW)/gridDim.x;
+// d_temp[int(x*tmp+(y*tmp)*imageW)] = 1;
+// d_temp[int(x*tmp+(tmp*(float(y)+0.5f)*imageW))] = 1;
+// d_temp[int(tmp*(float(x)+0.5f)+(y*tmp)*imageW)] = 1;
+// d_temp[int(tmp*(float(x)+0.5f)+(tmp*(float(y)+0.5f)*imageW))] = 1;
+ //if(gridDim.x==16) printf("hep %d %f\n",gridDim.x,t);
+ }
+ }
+ else {
+// float tmp= float(imageW)/gridDim.x;
+// d_temp[int(x*tmp+(y*tmp)*imageW)] = 1;
+// d_temp[int(x*tmp+(tmp*(float(y)+0.5f)*imageW))] = 1;
+// d_temp[int(tmp*(float(x)+0.5f)+(y*tmp)*imageW)] = 1;
+// d_temp[int(tmp*(float(x)+0.5f)+(tmp*(float(y)+0.5f)*imageW))] = 1;
+ //if(gridDim.x==16) printf("hep %d %f\n",gridDim.x,t);
+ }
+ //__syncthreads();
+}*/
+
+/*__global__ __device__ void rayCalc(float3 * A, float3 * u, float * prof, uint imageW, uint imageH, float df, float tPixel)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint id = x + y * imageW;
+
+ if( x < imageW && y < imageH )
+ {
+ matrice3x4 M(MView);
+ Rayon R;
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+ A[id] = R.A;
+ u[id] = R.u;
+ prof[id] = 1000.0f;
+ }
+}*/
+
+
+/*__global__ __device__ void rayTrace(uint * Obj, float * prof, float3 * A, float3 * u, uint imageW, uint imageH, float pas, float df, uint nObj)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint id = x + y * imageW;
+
+ if( x < imageW && y < imageH )
+ {
+ Sphere s(cnode[nObj].s);
+ float t;
+ s.C.x += pas;
+ Rayon R;
+ R.A = A[id];
+ R.u = u[id];
+ t = intersectionSphere(R,s.C,s.r);
+
+ if( t > 0.0f && t < prof[id] ) {
+ prof[id] = t;
+ Obj[id] = nObj;
+ }
+ }
+}*/
+/*
+__global__ __device__ void color(uint * result, uint * Obj, float * prof, float3 * A, float3 * u, uint imageW, uint imageH, float pas)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint id = x + y * imageW;
+
+ if( x < imageW && y < imageH )
+ {
+ float t(prof[id]);
+ if( t > 0.0f && t < 1000.0f ) {
+ Rayon R;
+ R.A = A[id];
+ R.u = u[id];
+ Sphere s(cnode[Obj[id]].s);
+ s.C.x += pas;
+ float4 f = make_float4(s.R,s.V,s.B,s.A)*(dot(getNormale(R.A+R.u*t,s.C),(-1.0f)*R.u));
+ result[id] = rgbaFloatToInt(f);
+ }
+ else {
+ result[id] = 0;
+ }
+ prof[id] = 100000.0f;
+ }
+}*/
+#ifdef DEBUG_RT_CUDA
+__device__ bool notShadowRay( float4* d_debug_float4, uint* d_debug_uint, int i, Node * node, float3 A, float3 u, float pas ) {
+
+#else
+__device__ bool notShadowRay( Node * node, float3 A, float3 u, float pas ) {
+#endif
+ float t(0.0f);
+ Node n;
+ Rayon ray;
+ float3 L(make_float3(10.0f,10.0f,10.0f)), tmp;
+ float dst(dot(tmp=(L-A),tmp));
+ ray.A = A+u*0.0001f;
+ ray.u = u;
+ for( int j(0); j < numObj && !t; j++ ) {
+ n = cnode[j];
+ n.s.C.x += pas;
+ if( n.fg ){
+ t = intersectionPlan(ray,n.s.C,n.s.C);
+ #ifdef DEVICE_EMU
+// printf("%d: j=%d, intersectionPlan t=%e\n", threadIdx.x, j, t);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ //d_debug_uint4[threadIdx.x*16+4*j+0]=10;
+// d_debug_float4[threadIdx.x*32+16*i+3*j+0].x = t;
+// d_debug_float4[threadIdx.x*32+16*i+3*j+0].y = 99999.9f;
+ #endif
+ }
+ else{
+#ifdef DEVICE_EMU
+ printf("%d: i=%d, n.s.C = %e %e %e\n", threadIdx.x, i, n.s.C.x, n.s.C.y, n.s.C.z);
+ printf("%d: i=%d, n.s.r = %e\n", threadIdx.x, i, n.s.r);
+#endif
+ #ifdef DEBUG_RT_CUDA
+ d_debug_float4[threadIdx.x*32+16*i+3*j+0].x = n.s.C.x;
+ d_debug_float4[threadIdx.x*32+16*i+3*j+0].y = n.s.C.y;
+ d_debug_float4[threadIdx.x*32+16*i+3*j+0].z = n.s.C.z;
+ d_debug_float4[threadIdx.x*32+16*i+3*j+0].w = n.s.r;
+ #endif
+ t = intersectionSphere(ray,n.s.C,n.s.r);
+ #ifdef DEVICE_EMU
+ printf("%d: j=%d, intersectionSphere t=%e\n", threadIdx.x, j, t);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ d_debug_float4[threadIdx.x*32+16*i+3*j+1].x = t;
+ d_debug_float4[threadIdx.x*32+16*i+3*j+1].y = 99999.9f;
+ #endif
+ }
+ if( t > 0.0f && dot(tmp=(A+u*t),tmp) > dst ){
+ t = 0.0f;
+ #ifdef DEVICE_EMU
+// printf("%d: j=%d, && dot t=%e\n", threadIdx.x, j, t);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+// d_debug_float4[threadIdx.x*32+16*i+3*j+2].x = t;
+// d_debug_float4[threadIdx.x*32+16*i+3*j+2].y = 99999.9f;
+ #endif
+ }
+ }
+ #ifdef DEVICE_EMU
+// printf("%d: t=%e\n", threadIdx.x, t);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+// d_debug_float4[threadIdx.x*32+16*i+13].x = t;
+// d_debug_float4[threadIdx.x*32+16*i+13].y = 99999.9f;
+ d_debug_float4[threadIdx.x*32+16*i+15].x = 88888.8f;
+ #endif
+ return t == 0.0f;
+}
+
+__device__ float float2int_pow20(float a)
+{
+ return a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a;
+}
+
+__device__ float float2int_pow50(float a)
+{
+ return a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a;
+
+}
+#ifdef DEBUG_RT_CUDA
+__global__ __device__ void render(float4* d_debug_float4, uint* d_debug_uint, uint * result, Node * dnode, uint imageW, uint imageH, float pas, float df)
+#else
+__global__ __device__ void render(uint * result, Node * dnode, uint imageW, uint imageH, float pas, float df)
+#endif
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint tid(__umul24(threadIdx.y, blockDim.x) + threadIdx.x);
+
+ uint id(x + y * imageW);
+ float4 pile[5];
+ uint Obj, nRec(5), n(0);
+ //__shared__ Node node[numObj];
+ float prof, tmp;
+
+ //if( tid < numObj ) node[tid] = cnode[tid];
+
+ for( int i(0); i < nRec; ++i )
+ pile[i] = make_float4(0.0f,0.0f,0.0f,1.0f);
+
+ if( x < imageW && y < imageH )
+ {
+ prof = 10000.0f;
+ result[id] = 0;
+ float tPixel(2.0f/float(min(imageW,imageH)));
+ float4 f(make_float4(0.0f,0.0f,0.0f,1.0f));
+ matrice3x4 M(MView);
+ Rayon R;
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+#ifdef DEVICE_EMU
+// printf("%d: R.A = %e %e %e\n", threadIdx.x, R.A.x, R.A.y, R.A.z);
+// printf("%d: R.u = %e %e %e\n", threadIdx.x, R.u.x, R.u.y, R.u.z);
+#endif
+#ifdef DEBUG_RT_CUDA
+// d_debug_float4[threadIdx.x*2+0].x= R.A.x;
+// d_debug_float4[threadIdx.x*2+0].y= R.A.y;
+// d_debug_float4[threadIdx.x*2+0].z= R.A.z;
+// d_debug_float4[threadIdx.x*2+1].x= R.u.x;
+// d_debug_float4[threadIdx.x*2+1].y= R.u.y;
+// d_debug_float4[threadIdx.x*2+1].z= R.u.z;
+#endif
+ __syncthreads();
+
+ for( int i(0); i < nRec && n == i; i++ ) {
+
+ for( int j(0); j < numObj; j++ ) {
+ Node nod(cnode[j]);
+ Sphere s(nod.s);
+ float t;
+ s.C.x += pas;
+ if( nod.fg )
+ t = intersectionPlan(R,s.C,s.C);
+ else
+ t = intersectionSphere(R,s.C,s.r);
+
+ if( t > 0.0f && t < prof ) {
+ prof = t;
+ Obj = j;
+ }
+ }
+#ifdef DEBUG_RT_CUDA
+ //d_debug_float4[threadIdx.x*5+i].x= prof;
+#endif
+#ifdef DEVICE_EMU
+// printf("%d: i=%d, t=%e\n", threadIdx.x, i, prof);
+#endif
+ float t = prof;
+ if( t > 0.0f && t < 10000.0f ) {
+ n++;
+ Node nod(cnode[Obj]);
+ Sphere s(nod.s);
+ s.C.x += pas;
+ float4 color(make_float4(s.R,s.V,s.B,s.A));
+ float3 P(R.A+R.u*t), L(normalize(make_float3(10.0f,10.0f,10.0f)-P)), V(normalize(R.A-P));
+ float3 N(nod.fg?getNormaleP(P):getNormale(P,s.C));
+ float3 Np(dot(V,N)<0.0f?(-1*N):N);
+ pile[i] = 0.05f * color;
+ #ifdef DEVICE_EMU
+// printf("%d: i=%d, pile[i] = %e %e %e %e\n", threadIdx.x, i, pile[i].x, pile[i].y, pile[i].z, pile[i].w);
+// printf("%d: i=%d, color = %e %e %e %e\n", threadIdx.x, i, color.x, color.y, color.z, color.w);
+// printf("%d: i=%d, P = %e %e %e\n", threadIdx.x, i, P.x, P.y, P.z);
+// printf("%d: i=%d, L = %e %e %e\n", threadIdx.x, i, L.x, L.y, L.z);
+// printf("%d: i=%d, V = %e %e %e\n", threadIdx.x, i, V.x, V.y, V.z);
+// printf("%d: i=%d, N = %e %e %e\n", threadIdx.x, i, N.x, N.y, N.z);
+// printf("%d: i=%d, Np = %e %e %e\n", threadIdx.x, i, Np.x, Np.y, Np.z);
+// printf("%d: i=%d, dot(Np,L) = %e\n", threadIdx.x, i, dot(Np,L));
+ //printf("%d: i=%d, notShadowRay(cnode,P,L,pas) = %d\n", threadIdx.x, i, (int) notShadowRay(cnode,P,L,pas));
+
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ //d_debug_float4[threadIdx.x*16+i*3+0]= pile[i];
+// d_debug_float4[threadIdx.x*16+i*8+0]= color;
+// d_debug_float4[threadIdx.x*16+i*8+1].x= P.x;d_debug_float4[threadIdx.x*16+i*8+1].y= P.y;d_debug_float4[threadIdx.x*16+i*8+1].z= P.z;
+// d_debug_float4[threadIdx.x*16+i*8+2].x= L.x;d_debug_float4[threadIdx.x*16+i*8+2].y= L.y;d_debug_float4[threadIdx.x*16+i*8+2].z= L.z;
+// d_debug_float4[threadIdx.x*16+i*8+3].x= V.x;d_debug_float4[threadIdx.x*16+i*8+3].y= V.y;d_debug_float4[threadIdx.x*16+i*8+3].z= V.z;
+// d_debug_float4[threadIdx.x*16+i*8+4].x= N.x;d_debug_float4[threadIdx.x*16+i*8+4].y= N.y;d_debug_float4[threadIdx.x*16+i*8+4].z= N.z;
+// d_debug_float4[threadIdx.x*16+i*8+5].x= Np.x;d_debug_float4[threadIdx.x*16+i*8+5].y= Np.y;d_debug_float4[threadIdx.x*16+i*8+5].z= Np.z;
+// d_debug_float4[threadIdx.x*16+i*8+6].x= dot(Np,L);
+ //d_debug_float4[threadIdx.x*16+i*8+7].x= (float) notShadowRay(cnode,P,L,pas);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ if( dot(Np,L) > 0.0f && notShadowRay(d_debug_float4, d_debug_uint, i, cnode,P,L,pas) ) {
+ #else
+ if( dot(Np,L) > 0.0f && notShadowRay(cnode,P,L,pas) ) {
+ #endif
+ //float3 Ri(2.0f*Np*dot(Np,L) - L);
+ float3 Ri(normalize(L+V));
+ //Ri = (L+V)/normalize(L+V);
+ pile[i] += 0.3f * color* (min(1.0f,dot(Np,L)));
+ #ifdef DEVICE_EMU
+// printf("%d: i=%d, pile[i] = %e %e %e %e\n", threadIdx.x, i, pile[i].x, pile[i].y, pile[i].z, pile[i].w);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ //d_debug_float4[threadIdx.x*16+i*3+1]= pile[i];
+ #endif
+ #ifdef FIXED_CONST_PARSE
+ tmp = 0.8f * pow(max(0.0f,min(1.0f,dot(Np,Ri))),50.0f);
+ #else
+ tmp = 0.8f * float2int_pow50(max(0.0f,min(1.0f,dot(Np,Ri))));
+ #endif
+ pile[i].x += tmp;
+ pile[i].y += tmp;
+ pile[i].z += tmp;
+ #ifdef DEVICE_EMU
+// printf("%d: i=%d, pile[i] = %e %e %e %e\n", threadIdx.x, i, pile[i].x, pile[i].y, pile[i].z, pile[i].w);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ //d_debug_float4[threadIdx.x*16+i*3+2]= pile[i];
+ #endif
+ }
+
+ R.u = 2.0f*N*dot(N,V) - V;
+ R.u = normalize(R.u);
+ R.A = P+R.u*0.0001f;
+ }
+ prof = 10000.0f;
+ }
+ #ifdef DEBUG_RT_CUDA
+ /*d_debug_float4[threadIdx.x*5+0]= pile[0];
+ d_debug_float4[threadIdx.x*5+1]= pile[1];
+ d_debug_float4[threadIdx.x*5+2]= pile[2];
+ d_debug_float4[threadIdx.x*5+3]= pile[3];
+ d_debug_float4[threadIdx.x*5+4]= pile[4];*/
+ #endif
+#ifdef DEVICE_EMU
+// printf("%d: pile[0] = %e %e %e %e\n", threadIdx.x, pile[0].x, pile[0].y, pile[0].z, pile[0].w);
+// printf("%d: pile[1] = %e %e %e %e\n", threadIdx.x, pile[1].x, pile[1].y, pile[1].z, pile[1].w);
+// printf("%d: pile[2] = %e %e %e %e\n", threadIdx.x, pile[2].x, pile[2].y, pile[2].z, pile[2].w);
+// printf("%d: pile[3] = %e %e %e %e\n", threadIdx.x, pile[3].x, pile[3].y, pile[3].z, pile[3].w);
+// printf("%d: pile[4] = %e %e %e %e\n", threadIdx.x, pile[4].x, pile[4].y, pile[4].z, pile[4].w);
+#endif
+ for( int i(n-1); i > 0; i-- )
+ pile[i-1] = pile[i-1] + 0.8f*pile[i];
+#ifdef DEVICE_EMU
+// printf("%d: pile[0] = %e %e %e %e\n", threadIdx.x, pile[0].x, pile[0].y, pile[0].z, pile[0].w);
+#endif
+ result[id] += rgbaFloatToInt(pile[0]);
+ }
+}
+
+/*__global__ __device__ void renderPixel(uint * result, Node * dnode, uint imageW, uint imageH, float pas, float df)
+{
+ uint id(blockIdx.x + __umul24(blockIdx.y, imageW));
+ uint tid(threadIdx.x), x(blockIdx.x), y(blockIdx.y);
+ Node node;
+ float t(0.0f), tPixel;
+ float4 Color(make_float4(0.0f,0.0f,0.0f,1.0f));
+ matrice3x4 M(MView);
+ Rayon R;
+ Sphere s;
+ __shared__ float T[numObj];
+ __shared__ uint Obj;
+
+ T[tid] = 10000.0f;
+
+ if( x < imageW && y < imageH && tid < numObj ) {
+ node = dnode[tid];
+ if( tid == 0 ) result[id] = 0;
+ tPixel = 2.0f/float(min(imageW,imageH));
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+
+ s = node.s;
+ s.C.x += pas;
+
+ if( node.fg )
+ t = intersectionPlan(R,s.C,s.C);
+ else
+ t = intersectionSphere(R,s.C,s.r);
+
+ T[tid] = t;
+
+ __syncthreads();
+
+ if( tid == 0 ) {
+ float tmp(t);
+ Obj = 0;
+ for( int i(1); i < numObj; i++ ) {
+ if( T[i] > 0.0f && ( tmp == 0.0f || T[i] < tmp ) ) {
+ tmp = T[i];
+ Obj = i;
+ }
+ }
+ }
+
+ __syncthreads();
+
+ if( tid == Obj && t > 0.0f ) {
+ s = node.s;
+ s.C.x += pas;
+ float3 P(R.A+R.u*t), L(normalize(make_float3(0,1,2)-P)), V(-1*R.u);
+ float3 N(node.fg?getNormaleP(P):getNormale(P,s.C));
+ if( dot(N,L) > 0.0f ) {
+ Color = 0.5f*make_float4(s.R,s.V,s.B,s.A)*(max(0.0f,dot(N,L)));
+ #ifdef FIXED_CONST_PARSE
+ Color += 0.8f*make_float4(1.0f,1.0f,1.0f,1.0f)*pow(max(0.0f,min(1.0f,dot(2.0f*N*dot(N,L)-L,V))),20.0f);
+ #else
+ Color += 0.8f*make_float4(1.0f,1.0f,1.0f,1.0f)*float2int_pow20(max(0.0f,min(1.0f,dot(2.0f*N*dot(N,L)-L,V))));
+ #endif
+ }
+ result[id] = rgbaFloatToInt(Color);
+ }
+ }
+
+}
+*/
+
+#endif // __RAYTRACING_KERNEL_H__
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 10:37:23 +0000