From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Thu, 15 Jul 2010 18:09:46 -0800 Subject: creating branch for adding support for CUDA 3.x and Fermi [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829] --- benchmarks/CUDA/RAY/EasyBMP.cpp | 1905 +++++++++++++++++++++ benchmarks/CUDA/RAY/EasyBMP.h | 86 + benchmarks/CUDA/RAY/EasyBMP_BMP.h | 86 + benchmarks/CUDA/RAY/EasyBMP_DataStructures.h | 104 ++ benchmarks/CUDA/RAY/EasyBMP_VariousBMPutilities.h | 43 + benchmarks/CUDA/RAY/Makefile | 50 + benchmarks/CUDA/RAY/README.GPGPU-Sim | 2 + benchmarks/CUDA/RAY/makebmp.cpp | 28 + benchmarks/CUDA/RAY/makebmp.h | 9 + benchmarks/CUDA/RAY/rayTracing.cu | 441 +++++ benchmarks/CUDA/RAY/rayTracing_kernel.cu | 658 +++++++ 11 files changed, 3412 insertions(+) create mode 100644 benchmarks/CUDA/RAY/EasyBMP.cpp create mode 100644 benchmarks/CUDA/RAY/EasyBMP.h create mode 100644 benchmarks/CUDA/RAY/EasyBMP_BMP.h create mode 100644 benchmarks/CUDA/RAY/EasyBMP_DataStructures.h create mode 100644 benchmarks/CUDA/RAY/EasyBMP_VariousBMPutilities.h create mode 100644 benchmarks/CUDA/RAY/Makefile create mode 100644 benchmarks/CUDA/RAY/README.GPGPU-Sim create mode 100644 benchmarks/CUDA/RAY/makebmp.cpp create mode 100644 benchmarks/CUDA/RAY/makebmp.h create mode 100644 benchmarks/CUDA/RAY/rayTracing.cu create mode 100644 benchmarks/CUDA/RAY/rayTracing_kernel.cu (limited to 'benchmarks/CUDA/RAY') 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: macklin01@users.sourceforge.net * +* 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 = 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 -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: macklin01@users.sourceforge.net * +* 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 +#include +#include +#include + +#ifndef EasyBMP +#define EasyBMP + +#ifdef __BCPLUSPLUS__ +// The Borland compiler must use this because something +// is wrong with their cstdio file. +#include +#else +#include +#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: macklin01@users.sourceforge.net * +* 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: macklin01@users.sourceforge.net * +* 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: macklin01@users.sourceforge.net * +* 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 +#include +#include +#include +#include "makebmp.h" + +/*#include +#include + +#include */ +#include + +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 + +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 + //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<<>>(d_debug_float4, d_debug_uint, d_output, d_node, width, height, anim, obs.getDistance()); + #else + render<<>>(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)); // + +} + +// 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<= 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 [-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 +__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__ -- cgit v1.3