From e276548b0927bc9f164a1d9c9e8705991cdac5d4 Mon Sep 17 00:00:00 2001 From: dbp Date: Thu, 31 Oct 2013 15:35:18 +0400 Subject: [PATCH] 0024130: Implementing ray tracing visualization core The purpose of this functionality is to bring a basic ray-tracing solution to existing OCCT visualization toolkit (TKOpenGL). Currently ray-tracing visualization core supports sharp shadows, specular reflections, transparency and adaptive anti-aliasing. However, the basis for all ray-tracing algorithms is versatile, allowing you to add new ray-tracing features easily (such as ambient occlusion). All ray-tracing computations are performed on the GPU using OpenCL framework, allowing real-time rendering performance. It is important to note, that real-time ray-tracing is possible using high-performance GPUs with support of OpenCL 1.1 and higher (such as NVIDIA GeForce 660 or ATI/AMD Radeon 7850). When using low-end GPUs (such as NVIDIA GeForce 640) the ray-tracing performance may slow down significantly. Therefore, even with NVIDIA GeForce 640 you can render scenes with the millions of triangles. The support of OpenCL-enabled CPUs and integrated graphics cards is not guaranteed. --- .gitattributes | 1 + adm/UDLIST | 1 + samples/qt/Common/res/antialiasing.png | Bin 0 -> 230 bytes samples/qt/Common/res/reflections.png | Bin 0 -> 217 bytes samples/qt/Common/res/shadows.png | Bin 0 -> 264 bytes samples/qt/Common/src/ApplicationCommon.cxx | 188 +- samples/qt/Common/src/ApplicationCommon.h | 16 +- samples/qt/Common/src/Common-icon.ts | 20 + samples/qt/Common/src/Common-string.ts | 42 +- samples/qt/Common/src/DocumentCommon.cxx | 57 +- samples/qt/Common/src/DocumentCommon.h | 13 +- samples/qt/Common/src/MDIWindow.cxx | 42 +- samples/qt/Common/src/MDIWindow.h | 22 +- samples/qt/Common/src/View.cxx | 85 +- samples/qt/Common/src/View.h | 11 +- samples/qt/IESample/IESample-vc10.sln | 2 +- src/Graphic3d/Graphic3d_CView.hxx | 22 +- src/OpenGl/EXTERNLIB | 1 + src/OpenGl/FILES | 8 + src/OpenGl/OpenGl_AABB.cxx | 137 ++ src/OpenGl/OpenGl_AABB.hxx | 88 + src/OpenGl/OpenGl_Caps.cxx | 3 +- src/OpenGl/OpenGl_Caps.hxx | 7 + src/OpenGl/OpenGl_Cl.hxx | 33 + src/OpenGl/OpenGl_Display_2.cxx | 2 +- src/OpenGl/OpenGl_GraphicDriver.cxx | 33 + src/OpenGl/OpenGl_GraphicDriver.hxx | 6 + src/OpenGl/OpenGl_GraphicDriver_7.cxx | 10 + src/OpenGl/OpenGl_Group.cxx | 70 +- src/OpenGl/OpenGl_Group.hxx | 42 +- src/OpenGl/OpenGl_LayerList.cxx | 24 + src/OpenGl/OpenGl_LayerList.hxx | 14 + src/OpenGl/OpenGl_PrimitiveArray.cxx | 6 +- src/OpenGl/OpenGl_PriorityList.hxx | 3 + src/OpenGl/OpenGl_RaytraceSource.cxx | 940 ++++++++ src/OpenGl/OpenGl_RaytraceTypes.hxx | 46 + src/OpenGl/OpenGl_SceneGeometry.cxx | 761 ++++++ src/OpenGl/OpenGl_SceneGeometry.hxx | 360 +++ src/OpenGl/OpenGl_Structure.cxx | 206 +- src/OpenGl/OpenGl_Structure.hxx | 61 +- src/OpenGl/OpenGl_View.cxx | 24 +- src/OpenGl/OpenGl_View.hxx | 26 +- src/OpenGl/OpenGl_View_2.cxx | 255 +- src/OpenGl/OpenGl_Workspace.cxx | 48 +- src/OpenGl/OpenGl_Workspace.hxx | 179 +- src/OpenGl/OpenGl_Workspace_Raytrace.cxx | 2176 ++++++++++++++++++ src/TKOpenGl/EXTERNLIB | 1 + src/V3d/V3d_View.cdl | 32 + src/V3d/V3d_View_5.cxx | 56 + src/ViewerTest/ViewerTest_ViewerCommands.cxx | 188 +- tests/bugs/parse.rules | 1 + tests/bugs/vis/bug24130 | 48 + 52 files changed, 6198 insertions(+), 219 deletions(-) create mode 100644 samples/qt/Common/res/antialiasing.png create mode 100644 samples/qt/Common/res/reflections.png create mode 100644 samples/qt/Common/res/shadows.png create mode 100644 src/OpenGl/OpenGl_AABB.cxx create mode 100644 src/OpenGl/OpenGl_AABB.hxx create mode 100644 src/OpenGl/OpenGl_Cl.hxx create mode 100644 src/OpenGl/OpenGl_RaytraceSource.cxx create mode 100644 src/OpenGl/OpenGl_RaytraceTypes.hxx create mode 100644 src/OpenGl/OpenGl_SceneGeometry.cxx create mode 100644 src/OpenGl/OpenGl_SceneGeometry.hxx create mode 100644 src/OpenGl/OpenGl_Workspace_Raytrace.cxx create mode 100644 tests/bugs/vis/bug24130 diff --git a/.gitattributes b/.gitattributes index e805ec9760..6f32bdc850 100644 --- a/.gitattributes +++ b/.gitattributes @@ -12,6 +12,7 @@ *.jxx eol=lf *.lxx eol=lf *.pxx eol=lf +*.cl eol=lf *.cdl eol=lf *.edl eol=lf *.yacc eol=lf diff --git a/adm/UDLIST b/adm/UDLIST index eb7824cb39..c6cea580ef 100644 --- a/adm/UDLIST +++ b/adm/UDLIST @@ -488,3 +488,4 @@ p BOPAlgo p BOPDS p BOPCol p BOPInt +r OpenCL diff --git a/samples/qt/Common/res/antialiasing.png b/samples/qt/Common/res/antialiasing.png new file mode 100644 index 0000000000000000000000000000000000000000..da8e5047a4ff2c2738718a56cfd47c1d19d24d1e GIT binary patch literal 230 zcmeAS@N?(olHy`uVBq!ia0vp^0wB!61|;P_|4#%`Ea{HEjtmSN`?>!lvI6;>1s;*b z3=DjSK$uZf!>a)(C{^MbQ4*Y=R#Ki=l*$m0n3-3i=jR%tP-d)Ws%L2E{@KYKs4B_R z#WBRfKl#W1|MtwP5^E$m99X#OzQ2<_94$HZzXDHEaxn8XZexZe9J0#N4SMWG6BaPu z;A!pTHDdl?DLP-qhNpm4Oy@}sOGTn}2R UAKHp91KP&m>FVdQ&MBb@07Pj*oB#j- literal 0 HcmV?d00001 diff --git a/samples/qt/Common/res/reflections.png b/samples/qt/Common/res/reflections.png new file mode 100644 index 0000000000000000000000000000000000000000..ad48a956150be0b6cc42f2a0b063865711017786 GIT binary patch literal 217 zcmeAS@N?(olHy`uVBq!ia0vp^0wB!61|;P_|4#%`Ea{HEjtmSN`?>!lvI6;>1s;*b z3=DjSL74G){)!Z!pj3%#L`iUdT1k0gQ7S_~VrE{6o}X)oLYc9ish**s`)4O}psG+$ z7sn6_|KuP4|JyUGO01jN*!a+1QzA{Xi}C$B(?9BL%{o%+*aAd4p17%6h`oCK`v2r+ z$0|)P7Y()vub6sD7>pkLoy_Z!IN_FK7K7%)hKH<0ml+wBY8~9U`j0vf&;kZeS3j3^ HP6!lvI6;>1s;*b z3=DjSL74G){)!Z!pj3%#L`iUdT1k0gQ7S_~VrE{6o}X)oLYc9ish**s`)4O}psE&6 z7sn6_|F@Gbavo3+aL)g7?je`o10`AKi}|X#e1A4{guY6*mus4%IAcx+M;fDo|GVu> z^>XVA-feF9E~j#=W6tJXkwpn>SqsmtJh7?4kgY|j #include +#include +#include + #include static ApplicationCommonWindow* stApp; @@ -56,13 +59,12 @@ void ApplicationCommonWindow::createStandardOperations() QString dir = getResourceDir() + QString( "/" ); - newIcon = QPixmap( dir+QObject::tr("ICON_NEW") ); - helpIcon = QPixmap( dir+QObject::tr("ICON_HELP") ); - closeIcon = QPixmap( dir+QObject::tr("ICON_CLOSE") ); + newIcon = QPixmap( dir + QObject::tr("ICON_NEW") ); + helpIcon = QPixmap( dir + QObject::tr("ICON_HELP") ); + closeIcon = QPixmap( dir + QObject::tr("ICON_CLOSE") ); - QAction * fileNewAction, * fileCloseAction, * fileQuitAction, - * viewToolAction, * viewStatusAction, - * helpAboutAction; + QAction * fileNewAction, * fileCloseAction, * filePrefUseVBOAction, + * fileQuitAction, * viewToolAction, * viewStatusAction, * helpAboutAction; fileNewAction = new QAction( newIcon, QObject::tr("MNU_NEW"), this ); fileNewAction->setToolTip( QObject::tr("TBR_NEW") ); @@ -78,6 +80,14 @@ void ApplicationCommonWindow::createStandardOperations() connect( fileCloseAction, SIGNAL( triggered() ) , this, SLOT( onCloseWindow() ) ); myStdActions.insert( FileCloseId, fileCloseAction ); + filePrefUseVBOAction = new QAction( QObject::tr("MNU_USE_VBO"), this ); + filePrefUseVBOAction->setToolTip( QObject::tr("TBR_USE_VBO") ); + filePrefUseVBOAction->setStatusTip( QObject::tr("TBR_USE_VBO") ); + filePrefUseVBOAction->setCheckable( true ); + filePrefUseVBOAction->setChecked( true ); + connect( filePrefUseVBOAction, SIGNAL( activated() ) , this, SLOT( onUseVBO() ) ); + myStdActions.insert( FilePrefUseVBOId, filePrefUseVBOAction ); + fileQuitAction = new QAction( QObject::tr("MNU_QUIT"), this ); fileQuitAction->setToolTip( QObject::tr("TBR_QUIT") ); fileQuitAction->setStatusTip( QObject::tr("TBR_QUIT") ); @@ -108,12 +118,18 @@ void ApplicationCommonWindow::createStandardOperations() connect( helpAboutAction, SIGNAL( triggered() ) , this, SLOT( onAbout() ) ); myStdActions.insert( HelpAboutId, helpAboutAction ); - // popuplate a menu with all actions + // create preferences menu + QMenu* aPrefMenu = new QMenu( QObject::tr("MNU_PREFERENCES") ); + aPrefMenu->addAction( filePrefUseVBOAction ); + + // popuplate a menu with all actions myFilePopup = new QMenu( this ); myFilePopup = menuBar()->addMenu( QObject::tr("MNU_FILE") ); myFilePopup->addAction( fileNewAction ); myFilePopup->addAction( fileCloseAction ); myFileSeparator = myFilePopup->addSeparator(); + myFilePopup->addMenu( aPrefMenu ); + myFileSeparator = myFilePopup->addSeparator(); myFilePopup->addAction( fileQuitAction ); // add a view menu @@ -124,7 +140,6 @@ void ApplicationCommonWindow::createStandardOperations() view->addAction( viewStatusAction ); // add a help menu - QMenu * help = new QMenu( this ); menuBar()->addSeparator(); help = menuBar()->addMenu( QObject::tr("MNU_HELP") ); @@ -136,6 +151,8 @@ void ApplicationCommonWindow::createStandardOperations() myStdToolBar->addAction( helpAboutAction ); myStdActions.at(FileCloseId)->setEnabled(myDocuments.count() > 0); + + myStdActions.at(FilePrefUseVBOId)->setEnabled( true ); } void ApplicationCommonWindow::createCasCadeOperations() @@ -181,9 +198,43 @@ void ApplicationCommonWindow::createCasCadeOperations() a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_DEL") ), QObject::tr("MNU_TOOL_DEL"), this ); a->setToolTip( QObject::tr("TBR_TOOL_DEL") ); a->setStatusTip( QObject::tr("TBR_TOOL_DEL") ); - connect( a, SIGNAL( triggered() ) , this, SLOT( onToolAction() ) ); + connect( a, SIGNAL( activated() ) , this, SLOT( onToolAction() ) ); myToolActions.insert( ToolDeleteId, a ); +#ifdef HAVE_OPENCL + + // populate a tool bar with some actions + myRaytraceBar = addToolBar( tr( "Ray-trace options" ) ); + + a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_SHADOWS") ), QObject::tr("MNU_TOOL_SHADOWS"), this ); + a->setToolTip( QObject::tr("TBR_TOOL_SHADOWS") ); + a->setStatusTip( QObject::tr("TBR_TOOL_SHADOWS") ); + a->setCheckable( true ); + a->setChecked( true ); + connect( a, SIGNAL( activated() ) , this, SLOT( onRaytraceAction() ) ); + myRaytraceActions.insert( ToolShadowsId, a ); + myRaytraceBar->addAction( a ); + + a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_REFLECTIONS") ), QObject::tr("MNU_TOOL_REFLECTIONS"), this ); + a->setToolTip( QObject::tr("TBR_TOOL_REFLECTIONS") ); + a->setStatusTip( QObject::tr("TBR_TOOL_REFLECTIONS") ); + a->setCheckable( true ); + a->setChecked( true ); + connect( a, SIGNAL( activated() ) , this, SLOT( onRaytraceAction() ) ); + myRaytraceActions.insert( ToolReflectionsId, a ); + myRaytraceBar->addAction( a ); + + a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_ANTIALIASING") ), QObject::tr("MNU_TOOL_ANTIALIASING"), this ); + a->setToolTip( QObject::tr("TBR_TOOL_ANTIALIASING") ); + a->setStatusTip( QObject::tr("TBR_TOOL_ANTIALIASING") ); + a->setCheckable( true ); + a->setChecked( false ); + connect( a, SIGNAL( activated() ) , this, SLOT( onRaytraceAction() ) ); + myRaytraceActions.insert( ToolAntialiasingId, a ); + myRaytraceBar->addAction( a ); + +#endif + QSignalMapper* sm = new QSignalMapper( this ); connect( sm, SIGNAL( mapped( int ) ), this, SLOT( onSetMaterial( int ) ) ); @@ -275,12 +326,22 @@ void ApplicationCommonWindow::windowsMenuAboutToShow() QString dir = getResourceDir() + QString( "/" ); - a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_NEW3D" ) ), QObject::tr( "MNU_WINDOW_NEW3D" ), this ); - a->setToolTip( QObject::tr( "TBR_WINDOW_NEW3D" ) ); - a->setStatusTip( QObject::tr( "TBR_WINDOW_NEW3D" ) ); + a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_NEW3D" ) ), QObject::tr( "MNU_WINDOW_NEW3D_GL" ), this ); + a->setToolTip( QObject::tr( "TBR_WINDOW_NEW3D_GL" ) ); + a->setStatusTip( QObject::tr( "TBR_WINDOW_NEW3D_GL" ) ); connect( a, SIGNAL( triggered() ), this, SLOT( onCreateNewView() ) ); myWindowPopup->addAction( a ); +#ifdef HAVE_OPENCL + + a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_NEW3D" ) ), QObject::tr( "MNU_WINDOW_NEW3D_RT" ), this ); + a->setToolTip( QObject::tr( "TBR_WINDOW_NEW3D_RT" ) ); + a->setStatusTip( QObject::tr( "TBR_WINDOW_NEW3D_RT" ) ); + connect( a, SIGNAL( activated() ), this, SLOT( onCreateNewViewRT() ) ); + myWindowPopup->addAction( a ); + +#endif + a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_CASCADE" ) ), QObject::tr( "MNU_WINDOW_CASCADE" ), this ); a->setToolTip( QObject::tr( "TBR_WINDOW_CASCADE" ) ); a->setStatusTip( QObject::tr( "TBR_WINDOW_CASCADE" ) ); @@ -401,7 +462,25 @@ DocumentCommon* ApplicationCommonWindow::onNewDoc() return aDoc; } -void ApplicationCommonWindow::onCloseWindow(){ +DocumentCommon* ApplicationCommonWindow::onNewDocRT() +{ + updateFileActions(); + DocumentCommon* aDoc = createNewDocument(); + aDoc->onCreateNewView(true); + onSelectionChanged(); + connect( aDoc, SIGNAL( sendCloseDocument( DocumentCommon* ) ), + this, SLOT( onCloseDocument( DocumentCommon* ) ) ); + connect( stWs, SIGNAL( windowActivated( QWidget* ) ), + this, SLOT( onWindowActivated( QWidget* ) ) ); + connect( aDoc, SIGNAL( selectionChanged() ), + this, SLOT( onSelectionChanged() ) ); + myDocuments.append( aDoc ); + myStdActions.at( FileCloseId )->setEnabled( myDocuments.count() > 0 ); + return aDoc; +} + +void ApplicationCommonWindow::onCloseWindow() +{ MDIWindow* m = (MDIWindow*)stWs->activeWindow(); if ( m ) { @@ -410,6 +489,27 @@ void ApplicationCommonWindow::onCloseWindow(){ } } +void ApplicationCommonWindow::onUseVBO() +{ + MDIWindow* w = ( MDIWindow* ) stWs->activeWindow(); + + if ( NULL == w ) + return; + + Handle(AIS_InteractiveContext) aContextAIS = w->getDocument()->getContext(); + + if (aContextAIS.IsNull()) + return; + + Handle(OpenGl_GraphicDriver) aDriver = + Handle(OpenGl_GraphicDriver)::DownCast (aContextAIS->CurrentViewer()->Driver()); + + if (!aDriver.IsNull()) + { + aDriver->ChangeOptions().vboDisable = Standard_True; + } +} + void ApplicationCommonWindow::onCloseDocument(DocumentCommon* theDoc) { myDocuments.removeAll( theDoc ); @@ -449,14 +549,38 @@ void ApplicationCommonWindow::onAbout() void ApplicationCommonWindow::onCreateNewView() { - DocumentCommon* doc = ((MDIWindow*) stWs->activeWindow())->getDocument(); - doc->onCreateNewView(); + MDIWindow* window = qobject_cast< MDIWindow* >( stWs->activeWindow() ); + window->getDocument()->onCreateNewView( false ); +} + +#ifdef HAVE_OPENCL + +void ApplicationCommonWindow::onCreateNewViewRT() +{ + MDIWindow* window = qobject_cast< MDIWindow* >( stWs->activeWindow() ); + window->getDocument()->onCreateNewView( true ); } +#endif + void ApplicationCommonWindow::onWindowActivated ( QWidget * w ) { - if ( w ) - ((MDIWindow*) w)->onWindowActivated(); + if (w == NULL) + { + return; + } + + MDIWindow* window = qobject_cast< MDIWindow* >(w); + + window->onWindowActivated(); + +#ifdef HAVE_OPENCL + + myRaytraceActions.at( ToolShadowsId )->setChecked (window->ShadowsEnabled()); + myRaytraceActions.at( ToolReflectionsId )->setChecked (window->ReflectionsEnabled()); + myRaytraceActions.at( ToolAntialiasingId )->setChecked (window->AntialiasingEnabled()); + +#endif } void ApplicationCommonWindow::onToolAction() @@ -484,6 +608,36 @@ void ApplicationCommonWindow::onToolAction() doc->onDelete(); } +#ifdef HAVE_OPENCL + +void ApplicationCommonWindow::onRaytraceAction() +{ + QAction* sentBy = (QAction*) sender(); + + DocumentCommon* doc = qobject_cast< MDIWindow* >( + ApplicationCommonWindow::getWorkspace()->activeWindow())->getDocument(); + + if( sentBy == myRaytraceActions.at( ToolShadowsId ) ) + { + bool flag = myRaytraceActions.at( ToolShadowsId )->isChecked(); + doc->onShadows( flag ); + } + + if( sentBy == myRaytraceActions.at( ToolReflectionsId ) ) + { + bool flag = myRaytraceActions.at( ToolReflectionsId )->isChecked(); + doc->onReflections( flag ); + } + + if( sentBy == myRaytraceActions.at( ToolAntialiasingId ) ) + { + bool flag = myRaytraceActions.at( ToolAntialiasingId )->isChecked(); + doc->onAntialiasing( flag ); + } +} + +#endif + void ApplicationCommonWindow::onSelectionChanged() { QWorkspace* ws = ApplicationCommonWindow::getWorkspace(); diff --git a/samples/qt/Common/src/ApplicationCommon.h b/samples/qt/Common/src/ApplicationCommon.h index 742af43012..787f458797 100755 --- a/samples/qt/Common/src/ApplicationCommon.h +++ b/samples/qt/Common/src/ApplicationCommon.h @@ -10,13 +10,15 @@ #include #include + class COMMONSAMPLE_EXPORT ApplicationCommonWindow: public QMainWindow { Q_OBJECT public: - enum { FileNewId, FileCloseId, FileQuitId, ViewToolId, ViewStatusId, HelpAboutId }; + enum { FileNewId, FilePrefUseVBOId, FileCloseId, FilePreferencesId, FileQuitId, ViewToolId, ViewStatusId, HelpAboutId }; enum { ToolWireframeId, ToolShadingId, ToolColorId, ToolMaterialId, ToolTransparencyId, ToolDeleteId }; + enum { ToolShadowsId, ToolReflectionsId, ToolAntialiasingId }; ApplicationCommonWindow(); ~ApplicationCommonWindow(); @@ -36,14 +38,22 @@ protected: public slots: DocumentCommon* onNewDoc(); + DocumentCommon* onNewDocRT(); void onCloseWindow(); + void onUseVBO(); virtual void onCloseDocument( DocumentCommon* theDoc ); virtual void onSelectionChanged(); virtual void onAbout(); - void onViewToolBar(); + void onViewToolBar(); void onViewStatusBar(); void onToolAction(); +#ifdef HAVE_OPENCL + void onRaytraceAction(); +#endif void onCreateNewView(); +#ifdef HAVE_OPENCL + void onCreateNewViewRT(); +#endif void onWindowActivated ( QWidget * w ); void windowsMenuAboutToShow(); void windowsMenuActivated( bool checked/*int id*/ ); @@ -67,11 +77,13 @@ private: QList myStdActions; QList myToolActions; + QList myRaytraceActions; QList myMaterialActions; QList myDocuments; QToolBar* myStdToolBar; QToolBar* myCasCadeBar; + QToolBar* myRaytraceBar; QMenu* myFilePopup; QMenu* myWindowPopup; QAction* myFileSeparator; diff --git a/samples/qt/Common/src/Common-icon.ts b/samples/qt/Common/src/Common-icon.ts index 806deda57f..a5be1b3cfc 100755 --- a/samples/qt/Common/src/Common-icon.ts +++ b/samples/qt/Common/src/Common-icon.ts @@ -86,6 +86,14 @@ ICON_NEW new.png + + ICON_NEW_GL + newGL.png + + + ICON_NEW_RT + newRT.png + ICON_VIEW_RIGHT view_right.png @@ -126,5 +134,17 @@ ICON_SAMPLE lamp.png + + ICON_TOOL_SHADOWS + shadows.png + + + ICON_TOOL_REFLECTIONS + reflections.png + + + ICON_TOOL_ANTIALIASING + antialiasing.png + diff --git a/samples/qt/Common/src/Common-string.ts b/samples/qt/Common/src/Common-string.ts index 86fff361a8..cd87a190d7 100755 --- a/samples/qt/Common/src/Common-string.ts +++ b/samples/qt/Common/src/Common-string.ts @@ -94,6 +94,14 @@ MNU_FILE &File + + MNU_PREFERENCES + &Preferences + + + MNU_USE_VBO + &Use VBO + MNU_GOLD &Gold @@ -130,6 +138,18 @@ MNU_TOOL_TRANS &Transpatency + + MNU_TOOL_SHADOWS + &Shadows + + + MNU_TOOL_REFLECTIONS + &Reflections + + + MNU_TOOL_ANTIALIASING + &Antialiasing + BTN_BRASS Brass @@ -234,6 +254,10 @@ MNU_CH_BACK &Change Background + + MNU_CH_ENV_MAP + &Environment Map + TBR_CH_BACK Change Background @@ -316,11 +340,23 @@ TBR_WINDOW_NEW3D - New 3d View + New 3D View + + + TBR_WINDOW_NEW3D_GL + New GL 3D View + + + MNU_WINDOW_NEW3D_GL + New GL 3D View + + + TBR_WINDOW_NEW3D_RT + &New RT 3D View - MNU_WINDOW_NEW3D - &New 3d View + MNU_WINDOW_NEW3D_RT + &New RT 3D View MNU_STATUS_BAR diff --git a/samples/qt/Common/src/DocumentCommon.cxx b/samples/qt/Common/src/DocumentCommon.cxx index d8de9df361..71647c8e50 100755 --- a/samples/qt/Common/src/DocumentCommon.cxx +++ b/samples/qt/Common/src/DocumentCommon.cxx @@ -50,12 +50,10 @@ myNbViews( 0 ) myViewer = Viewer( getenv("DISPLAY"), a3DName.ToExtString(), "", 1000.0, V3d_XposYnegZpos, Standard_True, Standard_True ); - myViewer->Init(); myViewer->SetDefaultLights(); myViewer->SetLightOn(); - myContext =new AIS_InteractiveContext(myViewer); - //onCreateNewView(); + myContext =new AIS_InteractiveContext(myViewer); } DocumentCommon::~DocumentCommon() @@ -67,15 +65,16 @@ ApplicationCommonWindow* DocumentCommon::getApplication() return myApp; } -MDIWindow* DocumentCommon::createNewMDIWindow() +MDIWindow* DocumentCommon::createNewMDIWindow( bool theRT ) { QWorkspace* ws = myApp->getWorkspace(); - return new MDIWindow( this, ws, 0 ); + return new MDIWindow( this, ws, 0, theRT ); } -void DocumentCommon::onCreateNewView() + +void DocumentCommon::onCreateNewView( bool theRT ) { QWorkspace* ws = myApp->getWorkspace(); - MDIWindow* w = createNewMDIWindow(); + MDIWindow* w = createNewMDIWindow( theRT ); if( !w ) return; @@ -228,4 +227,48 @@ void DocumentCommon::onDelete() getApplication()->onSelectionChanged(); } +#ifdef HAVE_OPENCL + +void DocumentCommon::onShadows( int state ) +{ + QWorkspace* ws = ApplicationCommonWindow::getWorkspace(); + + MDIWindow* window = qobject_cast< MDIWindow* >( ws->activeWindow() ); + + if( window == NULL ) + return; + + window->setRaytracedShadows( state ); + + myContext->UpdateCurrentViewer(); +} + +void DocumentCommon::onReflections( int state ) +{ + QWorkspace* ws = ApplicationCommonWindow::getWorkspace(); + + MDIWindow* window = qobject_cast< MDIWindow* >( ws->activeWindow() ); + + if( window == NULL ) + return; + window->setRaytracedReflections( state ); + + myContext->UpdateCurrentViewer(); +} + +void DocumentCommon::onAntialiasing( int state ) +{ + QWorkspace* ws = ApplicationCommonWindow::getWorkspace(); + + MDIWindow* window = qobject_cast< MDIWindow* >( ws->activeWindow() ); + + if( window == NULL ) + return; + + window->setRaytracedAntialiasing( state ); + + myContext->UpdateCurrentViewer(); +} + +#endif diff --git a/samples/qt/Common/src/DocumentCommon.h b/samples/qt/Common/src/DocumentCommon.h index 3dcefb4d59..b8347fc28a 100755 --- a/samples/qt/Common/src/DocumentCommon.h +++ b/samples/qt/Common/src/DocumentCommon.h @@ -28,7 +28,7 @@ public: void fitAll(); protected: - virtual MDIWindow* createNewMDIWindow(); + virtual MDIWindow* createNewMDIWindow( bool theRT = false ); signals: void selectionChanged(); @@ -36,10 +36,17 @@ signals: public slots: virtual void onCloseView( MDIWindow* ); - virtual void onCreateNewView(); + virtual void onCreateNewView( bool theRT = false ); virtual void onMaterial(); - virtual void onMaterial( int ); + virtual void onMaterial( int ); virtual void onDelete(); + +#ifdef HAVE_OPENCL + virtual void onShadows( int state ); + virtual void onReflections( int state ); + virtual void onAntialiasing( int state ); +#endif + void onWireframe(); void onShading(); void onColor(); diff --git a/samples/qt/Common/src/MDIWindow.cxx b/samples/qt/Common/src/MDIWindow.cxx index 86bc793575..d44da21fda 100755 --- a/samples/qt/Common/src/MDIWindow.cxx +++ b/samples/qt/Common/src/MDIWindow.cxx @@ -24,9 +24,17 @@ MDIWindow::MDIWindow(View* aView, myView = aView; myDocument = aDocument; + +#ifdef HAVE_OPENCL + + myShadowsEnabled = true; + myReflectionsEnabled = true; + myAntialiasingEnabled = false; + +#endif } -MDIWindow::MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags ) +MDIWindow::MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags, bool theRT ) : QMainWindow( parent, wflags ) { QFrame *vb = new QFrame( this ); @@ -39,7 +47,7 @@ MDIWindow::MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlag setCentralWidget( vb ); myDocument = aDocument; - myView = new View( myDocument->getContext(), vb ); + myView = new View( myDocument->getContext(), vb, theRT ); layout->addWidget( myView ); connect( myView, SIGNAL( selectionChanged() ), @@ -49,6 +57,14 @@ MDIWindow::MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlag resize( sizeHint() ); setFocusPolicy( Qt::StrongFocus ); + +#ifdef HAVE_OPENCL + + myShadowsEnabled = true; + myReflectionsEnabled = true; + myAntialiasingEnabled = false; + +#endif } MDIWindow::~MDIWindow() @@ -137,4 +153,26 @@ QSize MDIWindow::sizeHint() const return QSize( 450, 300 ); } +#ifdef HAVE_OPENCL + +void MDIWindow::setRaytracedShadows( int state ) +{ + myView->setRaytracedShadows( state ); + myShadowsEnabled = state; +} + +void MDIWindow::setRaytracedReflections( int state ) +{ + myView->setRaytracedReflections( state ); + myReflectionsEnabled = state; +} + +void MDIWindow::setRaytracedAntialiasing( int state ) +{ + myView->setRaytracedAntialiasing( state ); + myAntialiasingEnabled = state; +} + +#endif + diff --git a/samples/qt/Common/src/MDIWindow.h b/samples/qt/Common/src/MDIWindow.h index 59e7df8104..e2753ccc75 100755 --- a/samples/qt/Common/src/MDIWindow.h +++ b/samples/qt/Common/src/MDIWindow.h @@ -12,7 +12,7 @@ class COMMONSAMPLE_EXPORT MDIWindow: public QMainWindow Q_OBJECT public: - MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags ); + MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags, bool theRT = false ); MDIWindow( View* aView, DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags ); ~MDIWindow(); @@ -20,6 +20,18 @@ public: void fitAll(); virtual QSize sizeHint() const; +#ifdef HAVE_OPENCL + + void setRaytracedShadows( int state ); + void setRaytracedReflections( int state ); + void setRaytracedAntialiasing( int state ); + + bool ShadowsEnabled() { return myShadowsEnabled; } + bool ReflectionsEnabled() { return myReflectionsEnabled; } + bool AntialiasingEnabled() { return myAntialiasingEnabled; } + +#endif + signals: void selectionChanged(); void message(const QString&, int ); @@ -36,6 +48,14 @@ protected: protected: DocumentCommon* myDocument; View* myView; + +#ifdef HAVE_OPENCL + + bool myShadowsEnabled; + bool myReflectionsEnabled; + bool myAntialiasingEnabled; + +#endif }; #endif diff --git a/samples/qt/Common/src/View.cxx b/samples/qt/Common/src/View.cxx index a881a83543..920faceb68 100755 --- a/samples/qt/Common/src/View.cxx +++ b/samples/qt/Common/src/View.cxx @@ -11,12 +11,14 @@ #include #include #include +#include #include #include #include #include #include +#include #include #if defined(_WIN32) || defined(__WIN32__) @@ -52,9 +54,11 @@ static QCursor* globPanCursor = NULL; static QCursor* zoomCursor = NULL; static QCursor* rotCursor = NULL; -View::View( Handle(AIS_InteractiveContext) theContext, QWidget* parent ) +View::View( Handle(AIS_InteractiveContext) theContext, QWidget* parent, bool theRT ) : QWidget( parent ), -myViewActions( 0 ) +myIsRT( theRT ), +myViewActions( 0 ), +myBackMenu( NULL ) { #if !defined(_WIN32) && !defined(__WIN32__) && (!defined(__APPLE__) || defined(MACOSX_USE_GLX)) //XSynchronize( x11Display(),true ); // it is possible to use QApplication::syncX(); @@ -63,6 +67,9 @@ myViewActions( 0 ) myFirst = true; myContext = theContext; + //if (theRT) + // myContext->SetDisplayMode( AIS_DisplayMode::AIS_Shaded, 1 ); + myXmin = 0; myYmin = 0; myXmax = 0; @@ -160,6 +167,7 @@ myViewActions( 0 ) View::~View() { + delete myBackMenu; } void View::init() @@ -184,6 +192,9 @@ void View::init() } myView->SetBackgroundColor (Quantity_NOC_BLACK); myView->MustBeResized(); + + if (myIsRT) + myView->SetRaytracingMode(); } void View::paintEvent( QPaintEvent * ) @@ -299,6 +310,30 @@ void View::hlrOn() QApplication::restoreOverrideCursor(); } +void View::setRaytracedShadows( int state ) +{ + if ( state ) + myView->EnableRaytracedShadows(); + else + myView->DisableRaytracedShadows(); +} + +void View::setRaytracedReflections( int state ) +{ + if ( state ) + myView->EnableRaytracedReflections(); + else + myView->DisableRaytracedReflections(); +} + +void View::setRaytracedAntialiasing( int state ) +{ + if ( state ) + myView->EnableRaytracedAntialiasing(); + else + myView->DisableRaytracedAntialiasing(); +} + void View::updateToggled( bool isOn ) { QAction* sentBy = (QAction*)sender(); @@ -847,14 +882,26 @@ void View::Popup( const int /*x*/, const int /*y*/ ) } else { - QMenu* myBackMenu = new QMenu( 0 ); - QAction* a = new QAction( QObject::tr("MNU_CH_BACK"), this ); - a->setToolTip( QObject::tr("TBR_CH_BACK") ); - connect( a,SIGNAL( triggered() ), this, SLOT( onBackground() ) ); - myBackMenu->addAction( a ); - addItemInPopup(myBackMenu); + if (!myBackMenu) + { + myBackMenu = new QMenu( 0 ); + + QAction* a = new QAction( QObject::tr("MNU_CH_BACK"), this ); + a->setToolTip( QObject::tr("TBR_CH_BACK") ); + connect( a, SIGNAL( activated() ), this, SLOT( onBackground() ) ); + myBackMenu->addAction( a ); + addItemInPopup(myBackMenu); + + a = new QAction( QObject::tr("MNU_CH_ENV_MAP"), this ); + a->setToolTip( QObject::tr("TBR_CH_ENV_MAP") ); + connect( a, SIGNAL( activated() ), this, SLOT( onEnvironmentMap() ) ); + a->setCheckable( true ); + a->setChecked( false ); + myBackMenu->addAction( a ); + addItemInPopup(myBackMenu); + } + myBackMenu->exec( QCursor::pos() ); - delete myBackMenu; } if ( w ) w->setFocus(); @@ -946,6 +993,26 @@ void View::onBackground() myView->Redraw(); } +void View::onEnvironmentMap() +{ + if (myBackMenu->actions().at(1)->isChecked()) + { + QString fileName = QFileDialog::getOpenFileName(this, tr("Open File"), "", + tr("All Image Files (*.bmp *.gif *.jpg *.jpeg *.png *.tga)")); + + Handle(Graphic3d_TextureEnv) aTexture = new Graphic3d_TextureEnv( fileName.toAscii().data() ); + + myView->SetTextureEnv (aTexture); + myView->SetSurfaceDetail (V3d_TEX_ENVIRONMENT); + } + else + { + myView->SetSurfaceDetail (V3d_TEX_NONE); + } + + myView->Redraw(); +} + bool View::dump(Standard_CString theFile) { myView->Redraw(); diff --git a/samples/qt/Common/src/View.h b/samples/qt/Common/src/View.h index 50c4a7cdbb..dac41c39f2 100755 --- a/samples/qt/Common/src/View.h +++ b/samples/qt/Common/src/View.h @@ -26,7 +26,9 @@ public: ViewAxoId, ViewRotationId, ViewResetId, ViewHlrOffId, ViewHlrOnId }; View( Handle(AIS_InteractiveContext) theContext, - QWidget* parent); + QWidget* parent, + bool theRT = false ); + ~View(); virtual void init(); @@ -35,6 +37,10 @@ public: void noActiveActions(); bool isShadingMode(); + void setRaytracedShadows( int state ); + void setRaytracedReflections( int state ); + void setRaytracedAntialiasing( int state ); + static QString GetMessages( int type,TopAbs_ShapeEnum aSubShapeType, TopAbs_ShapeEnum aShapeType ); static QString GetShapeType( TopAbs_ShapeEnum aShapeType ); @@ -70,6 +76,7 @@ public slots: void hlrOff(); void updateToggled( bool ); void onBackground(); + void onEnvironmentMap(); protected: virtual void paintEvent( QPaintEvent* ); @@ -107,6 +114,7 @@ private: const int MaxX, const int MaxY, const bool Draw ); private: + bool myIsRT; bool myFirst; bool myDrawRect; // set when a rect is used for selection or magnify Handle(V3d_View) myView; @@ -119,6 +127,7 @@ private: Quantity_Factor myCurZoom; Standard_Boolean myHlrModeIsOn; QList* myViewActions; + QMenu* myBackMenu; QRubberBand* myRectBand; //!< selection rectangle rubber band }; diff --git a/samples/qt/IESample/IESample-vc10.sln b/samples/qt/IESample/IESample-vc10.sln index 70a166763d..0cad3679f0 100644 --- a/samples/qt/IESample/IESample-vc10.sln +++ b/samples/qt/IESample/IESample-vc10.sln @@ -16,4 +16,4 @@ Global GlobalSection(SolutionProperties) = preSolution HideSolutionNode = FALSE EndGlobalSection -EndGlobal +EndGlobal \ No newline at end of file diff --git a/src/Graphic3d/Graphic3d_CView.hxx b/src/Graphic3d/Graphic3d_CView.hxx index 93b524d7e1..f2b46ecc69 100755 --- a/src/Graphic3d/Graphic3d_CView.hxx +++ b/src/Graphic3d/Graphic3d_CView.hxx @@ -94,7 +94,12 @@ public: Backfacing (0), GDisplayCB (NULL), GClientData (NULL), - ptrFBO (NULL) + ptrFBO (NULL), + WasRedrawnGL (0), + IsRaytracing (0), + IsShadowsEnabled (1), + IsReflectionsEnabled (1), + IsAntialiasingEnabled (0) { // } @@ -131,6 +136,21 @@ public: void* ptrFBO; + //! Was the window redrawn by standard OpenGL? + mutable int WasRedrawnGL; + + //! Enables/disables OpenCL-based ray-tracing. + int IsRaytracing; + + //! Enables/disables ray-traced sharp shadows. + int IsShadowsEnabled; + + //! Enables/disables ray-traced reflections. + int IsReflectionsEnabled; + + //! Enables/disables ray-traced adaptive anti-aliasing. + int IsAntialiasingEnabled; + }; const Handle(Standard_Type)& TYPE(Graphic3d_CView); diff --git a/src/OpenGl/EXTERNLIB b/src/OpenGl/EXTERNLIB index 21580fc276..a586f17029 100755 --- a/src/OpenGl/EXTERNLIB +++ b/src/OpenGl/EXTERNLIB @@ -6,6 +6,7 @@ CSF_objc CSF_Appkit CSF_IOKit CSF_OpenGlLibs +CSF_OPENCL CSF_AviLibs CSF_FREETYPE CSF_GL2PS diff --git a/src/OpenGl/FILES b/src/OpenGl/FILES index 57ae623767..a1686f8f35 100755 --- a/src/OpenGl/FILES +++ b/src/OpenGl/FILES @@ -133,3 +133,11 @@ OpenGl_ShaderStates.cxx Handle_OpenGl_ShaderObject.hxx Handle_OpenGl_ShaderProgram.hxx Handle_OpenGl_ShaderManager.hxx +OpenGl_Cl.hxx +OpenGl_AABB.hxx +OpenGl_AABB.cxx +OpenGl_SceneGeometry.hxx +OpenGl_SceneGeometry.cxx +OpenGl_RaytraceTypes.hxx +OpenGl_RaytraceSource.cxx +OpenGl_Workspace_Raytrace.cxx \ No newline at end of file diff --git a/src/OpenGl/OpenGl_AABB.cxx b/src/OpenGl/OpenGl_AABB.cxx new file mode 100644 index 0000000000..6dcdc2b9ae --- /dev/null +++ b/src/OpenGl/OpenGl_AABB.cxx @@ -0,0 +1,137 @@ +// Created on: 2013-08-27 +// Created by: Denis BOGOLEPOV +// Copyright (c) 2013 OPEN CASCADE SAS +// +// The content of this file is subject to the Open CASCADE Technology Public +// License Version 6.5 (the "License"). You may not use the content of this file +// except in compliance with the License. Please obtain a copy of the License +// at http://www.opencascade.org and read it completely before using this file. +// +// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its +// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France. +// +// The Original Code and all software distributed under the License is +// distributed on an "AS IS" basis, without warranty of any kind, and the +// Initial Developer hereby disclaims all such warranties, including without +// limitation, any warranties of merchantability, fitness for a particular +// purpose or non-infringement. Please see the License for the specific terms +// and conditions governing the rights and limitations under the License. + + +#ifdef HAVE_CONFIG_H + #include +#endif + +#ifdef HAVE_OPENCL + +#include + +#include + + +OpenGl_AABB::OpenGl_AABB() : myMinPoint(), + myMaxPoint(), + myIsValid (false) +{ } + +OpenGl_AABB::OpenGl_AABB (const OpenGl_RTVec4f& thePoint) : myMinPoint (thePoint), + myMaxPoint (thePoint), + myIsValid (true) +{ } + +OpenGl_AABB::OpenGl_AABB (const OpenGl_RTVec4f& theMinPoint, + const OpenGl_RTVec4f& theMaxPoint) : myMinPoint (theMinPoint), + myMaxPoint (theMaxPoint), + myIsValid (true) +{ } + +OpenGl_AABB::OpenGl_AABB (const OpenGl_AABB& theVolume) : myMinPoint (theVolume.myMinPoint), + myMaxPoint (theVolume.myMaxPoint), + myIsValid (theVolume.myIsValid) +{ } + +void OpenGl_AABB::Add (const OpenGl_RTVec4f& thePoint) +{ + if (!myIsValid) + { + myMinPoint = thePoint; + myMaxPoint = thePoint; + myIsValid = true; + } + else + { + myMinPoint = OpenGl_RTVec4f (Min (myMinPoint.x(), thePoint.x()), + Min (myMinPoint.y(), thePoint.y()), + Min (myMinPoint.z(), thePoint.z()), + 1.f); + + myMaxPoint = OpenGl_RTVec4f (Max (myMaxPoint.x(), thePoint.x()), + Max (myMaxPoint.y(), thePoint.y()), + Max (myMaxPoint.z(), thePoint.z()), + 1.f); + } +} + +void OpenGl_AABB::Combine (const OpenGl_AABB& theVolume) +{ + if (!theVolume.myIsValid) + return; + + if (!myIsValid) + { + myMinPoint = theVolume.myMinPoint; + myMaxPoint = theVolume.myMaxPoint; + myIsValid = true; + } + else + { + myMinPoint = OpenGl_RTVec4f (Min (myMinPoint.x(), theVolume.myMinPoint.x()), + Min (myMinPoint.y(), theVolume.myMinPoint.y()), + Min (myMinPoint.z(), theVolume.myMinPoint.z()), + 1.f); + + myMaxPoint = OpenGl_RTVec4f (Max (myMaxPoint.x(), theVolume.myMaxPoint.x()), + Max (myMaxPoint.y(), theVolume.myMaxPoint.y()), + Max (myMaxPoint.z(), theVolume.myMaxPoint.z()), + 1.f); + } +} + +OpenGl_AABB OpenGl_AABB::Added (const OpenGl_RTVec4f& thePoint) const +{ + OpenGl_AABB result (*this); + + result.Add (thePoint); + + return result; +} + +OpenGl_AABB OpenGl_AABB::Combined (const OpenGl_AABB& theVolume) const +{ + OpenGl_AABB result (*this); + + result.Combine (theVolume); + + return result; +} + +void OpenGl_AABB::Clear() +{ + myIsValid = false; +} + +OpenGl_RTVec4f OpenGl_AABB::Size() const +{ + return myMaxPoint - myMinPoint; +} + +float OpenGl_AABB::Area() const +{ + const float aXLen = myMaxPoint.x() - myMinPoint.x(); + const float aYLen = myMaxPoint.y() - myMinPoint.y(); + const float aZLen = myMaxPoint.z() - myMinPoint.z(); + + return ( aXLen * aYLen + aXLen * aZLen + aZLen * aYLen ) * 2.f; +} + +#endif \ No newline at end of file diff --git a/src/OpenGl/OpenGl_AABB.hxx b/src/OpenGl/OpenGl_AABB.hxx new file mode 100644 index 0000000000..9893406ad7 --- /dev/null +++ b/src/OpenGl/OpenGl_AABB.hxx @@ -0,0 +1,88 @@ +// Created on: 2013-08-27 +// Created by: Denis BOGOLEPOV +// Copyright (c) 2013 OPEN CASCADE SAS +// +// The content of this file is subject to the Open CASCADE Technology Public +// License Version 6.5 (the "License"). You may not use the content of this file +// except in compliance with the License. Please obtain a copy of the License +// at http://www.opencascade.org and read it completely before using this file. +// +// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its +// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France. +// +// The Original Code and all software distributed under the License is +// distributed on an "AS IS" basis, without warranty of any kind, and the +// Initial Developer hereby disclaims all such warranties, including without +// limitation, any warranties of merchantability, fitness for a particular +// purpose or non-infringement. Please see the License for the specific terms +// and conditions governing the rights and limitations under the License. + + +#ifndef _OpenGl_AABB_Header +#define _OpenGl_AABB_Header + +#include + + +//! Axis aligned bounding box (AABB). +class OpenGl_AABB +{ +public: + + //! Creates default (invalid) bounding volume. + OpenGl_AABB(); + + //! Creates bounding volume of given point. + OpenGl_AABB (const OpenGl_RTVec4f& thePoint); + + //! Creates copy of another bounding volume. + OpenGl_AABB (const OpenGl_AABB& theVolume); + + //! Creates bounding volume from min and max points. + OpenGl_AABB (const OpenGl_RTVec4f& theMinPoint, + const OpenGl_RTVec4f& theMaxPoint); + + //! Is object represents uninitialized volume? + bool IsVoid() const { return !myIsValid; } + + //! Appends new point to the volume. + void Add (const OpenGl_RTVec4f& theVector); + //! Combines the volume with another volume. + void Combine (const OpenGl_AABB& theVolume); + + //! Returns new volume created by appending a point to current volume. + OpenGl_AABB Added (const OpenGl_RTVec4f& thePoint) const; + //! Returns new volume created by combining with specified volume. + OpenGl_AABB Combined (const OpenGl_AABB& theVolume) const; + + //! Clears bounding volume (makes object invalid). + void Clear(); + + //! Evaluates surface area of bounding volume. + float Area() const; + + //! Return diagonal of bounding volume. + OpenGl_RTVec4f Size() const; + + //! Returns minimum point of bounding volume. + const OpenGl_RTVec4f& CornerMin() const { return myMinPoint; } + //! Returns maximum point of bounding volume. + const OpenGl_RTVec4f& CornerMax() const { return myMaxPoint; } + + //! Returns minimum point of bounding volume. + OpenGl_RTVec4f& CornerMin() { return myMinPoint; } + //! Returns maximum point of bounding volume. + OpenGl_RTVec4f& CornerMax() { return myMaxPoint; } + +private: + + //! Minimum point of bounding volume. + OpenGl_RTVec4f myMinPoint; + //! Maximum point of bounding volume. + OpenGl_RTVec4f myMaxPoint; + + //! Is bounding volume valid (up to date)? + bool myIsValid; +}; + +#endif diff --git a/src/OpenGl/OpenGl_Caps.cxx b/src/OpenGl/OpenGl_Caps.cxx index 39c14a818a..5b57f7b218 100644 --- a/src/OpenGl/OpenGl_Caps.cxx +++ b/src/OpenGl/OpenGl_Caps.cxx @@ -35,7 +35,8 @@ OpenGl_Caps::OpenGl_Caps() #else contextDebug (Standard_False), #endif - contextNoAccel (Standard_False) + contextNoAccel (Standard_False), + keepArrayData (Standard_False) { // } diff --git a/src/OpenGl/OpenGl_Caps.hxx b/src/OpenGl/OpenGl_Caps.hxx index 4ca46fc4cb..a7935bf3d3 100644 --- a/src/OpenGl/OpenGl_Caps.hxx +++ b/src/OpenGl/OpenGl_Caps.hxx @@ -67,6 +67,13 @@ public: //! @name context creation parameters */ Standard_Boolean contextNoAccel; + /** + * Disables freeing CPU memory after building VBOs. + * + * OFF by default. + */ + Standard_Boolean keepArrayData; + public: //! @name class methods //! Default constructor - initialize with most optimal values. diff --git a/src/OpenGl/OpenGl_Cl.hxx b/src/OpenGl/OpenGl_Cl.hxx new file mode 100644 index 0000000000..2bfcf33fd6 --- /dev/null +++ b/src/OpenGl/OpenGl_Cl.hxx @@ -0,0 +1,33 @@ +// Created on: 2013-10-15 +// Created by: Denis BOGOLEPOV +// Copyright (c) 2013 OPEN CASCADE SAS +// +// The content of this file is subject to the Open CASCADE Technology Public +// License Version 6.5 (the "License"). You may not use the content of this file +// except in compliance with the License. Please obtain a copy of the License +// at http://www.opencascade.org and read it completely before using this file. +// +// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its +// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France. +// +// The Original Code and all software distributed under the License is +// distributed on an "AS IS" basis, without warranty of any kind, and the +// Initial Developer hereby disclaims all such warranties, including without +// limitation, any warranties of merchantability, fitness for a particular +// purpose or non-infringement. Please see the License for the specific terms +// and conditions governing the rights and limitations under the License. + +#ifndef _OpenGl_Cl_H__ +#define _OpenGl_Cl_H__ + +// cl_gl.h includes OpenGL headers - make sure our stuff is included in right order +#include + +#if defined(__APPLE__) || defined(__MACOSX) + #include +#else + #include + #include +#endif + +#endif // _OpenGl_Cl_H__ diff --git a/src/OpenGl/OpenGl_Display_2.cxx b/src/OpenGl/OpenGl_Display_2.cxx index 5b04b88f92..5465b52c27 100644 --- a/src/OpenGl/OpenGl_Display_2.cxx +++ b/src/OpenGl/OpenGl_Display_2.cxx @@ -23,7 +23,7 @@ #include #ifdef HAVE_CONFIG_H -# include + #include #endif #ifdef HAVE_GL2PS diff --git a/src/OpenGl/OpenGl_GraphicDriver.cxx b/src/OpenGl/OpenGl_GraphicDriver.cxx index be8fcd7b01..18ca913a3d 100755 --- a/src/OpenGl/OpenGl_GraphicDriver.cxx +++ b/src/OpenGl/OpenGl_GraphicDriver.cxx @@ -17,6 +17,11 @@ // purpose or non-infringement. Please see the License for the specific terms // and conditions governing the rights and limitations under the License. + +#ifdef HAVE_CONFIG_H + #include +#endif + #include #include @@ -154,6 +159,34 @@ Standard_Boolean OpenGl_GraphicDriver::SetImmediateModeDrawToFront (const Graphi return Standard_False; } +// ======================================================================= +// function : GetOpenClDeviceInfo +// purpose : Returns information about device used for computations +// ======================================================================= +#ifndef HAVE_OPENCL + +Standard_Boolean OpenGl_GraphicDriver::GetOpenClDeviceInfo (const Graphic3d_CView&, + NCollection_DataMap&) +{ + return Standard_True; +} + +#else + +Standard_Boolean OpenGl_GraphicDriver::GetOpenClDeviceInfo (const Graphic3d_CView& theCView, + NCollection_DataMap& theInfo) +{ + + if (theCView.ViewId == -1 || theCView.ptrView == NULL) + { + return Standard_False; + } + + return reinterpret_cast (theCView.ptrView)->WS->GetOpenClDeviceInfo (theInfo); +} + +#endif + // ======================================================================= // function : BeginAddMode // purpose : diff --git a/src/OpenGl/OpenGl_GraphicDriver.hxx b/src/OpenGl/OpenGl_GraphicDriver.hxx index 39415f6c5f..a09ad544e9 100644 --- a/src/OpenGl/OpenGl_GraphicDriver.hxx +++ b/src/OpenGl/OpenGl_GraphicDriver.hxx @@ -319,6 +319,12 @@ public: //! Method to setup UserDraw callback Standard_EXPORT OpenGl_UserDrawCallback_t& UserDrawCallback(); +public: + + //! Returns information about OpenCL device used for computations. + Standard_EXPORT Standard_Boolean GetOpenClDeviceInfo (const Graphic3d_CView& theCView, + NCollection_DataMap& theInfo); + private: //! Method to retrieve valid GL context. diff --git a/src/OpenGl/OpenGl_GraphicDriver_7.cxx b/src/OpenGl/OpenGl_GraphicDriver_7.cxx index c1d9f57469..68d3aaa186 100755 --- a/src/OpenGl/OpenGl_GraphicDriver_7.cxx +++ b/src/OpenGl/OpenGl_GraphicDriver_7.cxx @@ -203,6 +203,16 @@ void OpenGl_GraphicDriver::Redraw (const Graphic3d_CView& ACView, const Standard_Integer /*width*/, const Standard_Integer /*height*/) { + if (!myCaps->vboDisable && ACView.IsRaytracing) + { + if (ACView.WasRedrawnGL) + { + myDeviceLostFlag = Standard_True; + } + + myCaps->keepArrayData = Standard_True; + } + const OpenGl_CView *aCView = (const OpenGl_CView *)ACView.ptrView; if (aCView) { diff --git a/src/OpenGl/OpenGl_Group.cxx b/src/OpenGl/OpenGl_Group.cxx index 5c6a63ce4d..b1151fb306 100644 --- a/src/OpenGl/OpenGl_Group.cxx +++ b/src/OpenGl/OpenGl_Group.cxx @@ -17,23 +17,37 @@ // purpose or non-infringement. Please see the License for the specific terms // and conditions governing the rights and limitations under the License. -#include +#ifdef HAVE_CONFIG_H + #include +#endif + +#include #include +#include #include // ======================================================================= // function : OpenGl_Group // purpose : // ======================================================================= -OpenGl_Group::OpenGl_Group () -: myAspectLine (NULL), - myAspectFace (NULL), - myAspectMarker (NULL), - myAspectText (NULL), - myFirst (NULL), - myLast (NULL) +#ifndef HAVE_OPENCL +OpenGl_Group::OpenGl_Group() +#else +OpenGl_Group::OpenGl_Group (const OpenGl_Structure* theAncestorStructure) +#endif +: myAspectLine(NULL), + myAspectFace(NULL), + myAspectMarker(NULL), + myAspectText(NULL), + myFirst(NULL), + myLast(NULL) { +#ifdef HAVE_OPENCL + myAncestorStructure = theAncestorStructure; + myIsRaytracable = Standard_False; + myModificationState = 0; // initial state +#endif } // ======================================================================= @@ -89,6 +103,18 @@ void OpenGl_Group::SetAspectFace (const CALL_DEF_CONTEXTFILLAREA& theAspect, anAspectFace->SetAspect (theAspect); AddElement (TelNil/*TelAspectFace*/, anAspectFace); } + +#ifdef HAVE_OPENCL + if (myIsRaytracable) + { + myModificationState++; + + if (myAncestorStructure != NULL) + { + myAncestorStructure->UpdateStateWithAncestorStructures(); + } + } +#endif } // ======================================================================= @@ -141,15 +167,29 @@ void OpenGl_Group::SetAspectText (const CALL_DEF_CONTEXTTEXT& theAspect, // function : AddElement // purpose : // ======================================================================= -void OpenGl_Group::AddElement (const TelType AType, OpenGl_Element *AElem ) +void OpenGl_Group::AddElement (const TelType theType, OpenGl_Element *theElem) { - OpenGl_ElementNode *node = new OpenGl_ElementNode(); + OpenGl_ElementNode *aNode = new OpenGl_ElementNode(); + + aNode->type = theType; + aNode->elem = theElem; + aNode->next = NULL; + (myLast? myLast->next : myFirst) = aNode; + myLast = aNode; + +#ifdef HAVE_OPENCL + if (OpenGl_Raytrace::IsRaytracedElement (aNode)) + { + myModificationState++; + myIsRaytracable = Standard_True; - node->type = AType; - node->elem = AElem; - node->next = NULL; - (myLast? myLast->next : myFirst) = node; - myLast = node; + if (myAncestorStructure != NULL) + { + myAncestorStructure->UpdateStateWithAncestorStructures(); + myAncestorStructure->SetRaytracableWithAncestorStructures(); + } + } +#endif } // ======================================================================= diff --git a/src/OpenGl/OpenGl_Group.hxx b/src/OpenGl/OpenGl_Group.hxx index 6ed3e996de..85c085c17c 100644 --- a/src/OpenGl/OpenGl_Group.hxx +++ b/src/OpenGl/OpenGl_Group.hxx @@ -33,6 +33,7 @@ #include class OpenGl_Group; +class OpenGl_Structure; typedef NCollection_List OpenGl_ListOfGroup; @@ -46,10 +47,13 @@ struct OpenGl_ElementNode class OpenGl_Group : public OpenGl_Element { - public: +#ifndef HAVE_OPENCL OpenGl_Group(); +#else + OpenGl_Group (const OpenGl_Structure* theAncestorStructure); +#endif void SetAspectLine (const CALL_DEF_CONTEXTLINE& theAspect, const Standard_Boolean IsGlobal = Standard_True); void SetAspectFace (const CALL_DEF_CONTEXTFILLAREA& theAspect, const Standard_Boolean IsGlobal = Standard_True); @@ -61,19 +65,41 @@ public: virtual void Render (const Handle(OpenGl_Workspace)& theWorkspace) const; virtual void Release (const Handle(OpenGl_Context)& theGlCtx); + //! Returns first OpenGL element node of the group. + const OpenGl_ElementNode* FirstNode() const { return myFirst; } + + //! Returns OpenGL face aspect. + const OpenGl_AspectFace* AspectFace() const { return myAspectFace; } + +#ifdef HAVE_OPENCL + + //! Returns modification state for ray-tracing. + Standard_Size ModificationState() const { return myModificationState; } + + //! Is the group ray-tracable (contains ray-tracable elements)? + Standard_Boolean IsRaytracable() const { return myIsRaytracable; } + +#endif + protected: virtual ~OpenGl_Group(); protected: - OpenGl_AspectLine* myAspectLine; - OpenGl_AspectFace* myAspectFace; - OpenGl_AspectMarker* myAspectMarker; - OpenGl_AspectText* myAspectText; - - OpenGl_ElementNode* myFirst; - OpenGl_ElementNode* myLast; + OpenGl_AspectLine* myAspectLine; + OpenGl_AspectFace* myAspectFace; + OpenGl_AspectMarker* myAspectMarker; + OpenGl_AspectText* myAspectText; + + OpenGl_ElementNode* myFirst; + OpenGl_ElementNode* myLast; + +#ifdef HAVE_OPENCL + const OpenGl_Structure* myAncestorStructure; + Standard_Boolean myIsRaytracable; + Standard_Size myModificationState; +#endif public: diff --git a/src/OpenGl/OpenGl_LayerList.cxx b/src/OpenGl/OpenGl_LayerList.cxx index cc25f07493..2e4b9923f3 100644 --- a/src/OpenGl/OpenGl_LayerList.cxx +++ b/src/OpenGl/OpenGl_LayerList.cxx @@ -19,6 +19,10 @@ // and conditions governing the rights and limitations under the License. +#ifdef HAVE_CONFIG_H + #include +#endif + #include #include @@ -152,6 +156,10 @@ void OpenGl_LayerList::AddStructure (const OpenGl_Structure *theStructure, aList.Add (theStructure, thePriority); myNbStructures++; + + // Note: In ray-tracing mode we don't modify modification + // state here. It is redundant, because the possible changes + // will be handled in the loop for structures } //======================================================================= @@ -173,6 +181,14 @@ void OpenGl_LayerList::RemoveStructure (const OpenGl_Structure *theStructure, if (aList.Remove (theStructure) >= 0) { myNbStructures--; + +#ifdef HAVE_OPENCL + if (theStructure->IsRaytracable()) + { + myModificationState++; + } +#endif + return; } @@ -188,6 +204,14 @@ void OpenGl_LayerList::RemoveStructure (const OpenGl_Structure *theStructure, if (aScanList.Remove (theStructure) >= 0) { myNbStructures--; + +#ifdef HAVE_OPENCL + if (theStructure->IsRaytracable()) + { + myModificationState++; + } +#endif + return; } } diff --git a/src/OpenGl/OpenGl_LayerList.hxx b/src/OpenGl/OpenGl_LayerList.hxx index eceae346cc..73828ec8ff 100644 --- a/src/OpenGl/OpenGl_LayerList.hxx +++ b/src/OpenGl/OpenGl_LayerList.hxx @@ -80,7 +80,17 @@ class OpenGl_LayerList //! Render this element void Render (const Handle(OpenGl_Workspace) &theWorkspace) const; + + //! Returns the set of OpenGL Z-layers. + const OpenGl_SequenceOfLayers& Layers() const { return myLayers; } +#ifdef HAVE_OPENCL + + //! Returns structure modification state (for ray-tracing). + Standard_Size ModificationState() const { return myModificationState; } + +#endif + private: //! Get default layer @@ -94,6 +104,10 @@ class OpenGl_LayerList Standard_Integer myNbPriorities; Standard_Integer myNbStructures; +#ifdef HAVE_OPENCL + mutable Standard_Size myModificationState; +#endif + public: DEFINE_STANDARD_ALLOC }; diff --git a/src/OpenGl/OpenGl_PrimitiveArray.cxx b/src/OpenGl/OpenGl_PrimitiveArray.cxx index ceda9b74dd..0a5b0b2fc4 100755 --- a/src/OpenGl/OpenGl_PrimitiveArray.cxx +++ b/src/OpenGl/OpenGl_PrimitiveArray.cxx @@ -149,7 +149,11 @@ Standard_Boolean OpenGl_PrimitiveArray::BuildVBO (const Handle(OpenGl_Workspace) } } - clearMemoryOwn(); + if (!aGlCtx->caps->keepArrayData) + { + clearMemoryOwn(); + } + return Standard_True; } diff --git a/src/OpenGl/OpenGl_PriorityList.hxx b/src/OpenGl/OpenGl_PriorityList.hxx index 8ea81609cf..7e43b7c06a 100644 --- a/src/OpenGl/OpenGl_PriorityList.hxx +++ b/src/OpenGl/OpenGl_PriorityList.hxx @@ -58,6 +58,9 @@ class OpenGl_PriorityList //! or less). Returns Standard_False if the list can not be accepted. Standard_Boolean Append (const OpenGl_PriorityList& theOther); + //! Returns array of OpenGL structures. + const OpenGl_ArrayOfStructure& ArrayOfStructures() const { return myArray; } + protected: OpenGl_ArrayOfStructure myArray; diff --git a/src/OpenGl/OpenGl_RaytraceSource.cxx b/src/OpenGl/OpenGl_RaytraceSource.cxx new file mode 100644 index 0000000000..baf9104b06 --- /dev/null +++ b/src/OpenGl/OpenGl_RaytraceSource.cxx @@ -0,0 +1,940 @@ +// Created on: 2013-10-16 +// Created by: Denis BOGOLEPOV +// Copyright (c) 2013 OPEN CASCADE SAS +// +// The content of this file is subject to the Open CASCADE Technology Public +// License Version 6.5 (the "License"). You may not use the content of this file +// except in compliance with the License. Please obtain a copy of the License +// at http://www.opencascade.org and read it completely before using this file. +// +// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its +// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France. +// +// The Original Code and all software distributed under the License is +// distributed on an "AS IS" basis, without warranty of any kind, and the +// Initial Developer hereby disclaims all such warranties, including without +// limitation, any warranties of merchantability, fitness for a particular +// purpose or non-infringement. Please see the License for the specific terms +// and conditions governing the rights and limitations under the License. + +#ifdef HAVE_CONFIG_H + #include +#endif + +#ifdef HAVE_OPENCL + +#define EOL "\n" + +extern const char THE_RAY_TRACE_OPENCL_SOURCE[] = + + ///////////////////////////////////////////////////////////////////////////////////////// + // Specific data types + EOL + //! Stores ray parameters. + EOL" typedef struct __SRay" + EOL" {" + EOL" float4 Origin;" + EOL" float4 Direct;" + EOL" }" + EOL" SRay;" + EOL + //! Stores parameters of intersection point. + EOL" typedef struct __SIntersect" + EOL" {" + EOL" float4 Normal;" + EOL" float Time;" + EOL" float U;" + EOL" float V;" + EOL" }" + EOL" SIntersect;" + EOL + EOL + ///////////////////////////////////////////////////////////////////////////////////////// + // Some useful vector constants + EOL + EOL" #define ZERO ( float4 )( 0.f, 0.f, 0.f, 0.f )" + EOL" #define UNIT ( float4 )( 1.f, 1.f, 1.f, 0.f )" + EOL + EOL" #define AXIS_X ( float4 )( 1.f, 0.f, 0.f, 0.f )" + EOL" #define AXIS_Y ( float4 )( 0.f, 1.f, 0.f, 0.f )" + EOL" #define AXIS_Z ( float4 )( 0.f, 0.f, 1.f, 0.f )" + EOL + EOL + ///////////////////////////////////////////////////////////////////////////////////////// + // Support functions + EOL + // ======================================================================= + // function : GenerateRay + // purpose : Generates primary ray for current work item + // ======================================================================= + EOL" void GenerateRay (SRay* theRay," + EOL" const float theX," + EOL" const float theY," + EOL" const int theSizeX," + EOL" const int theSizeY," + EOL" const float16 theOrigins," + EOL" const float16 theDirects)" + EOL" {" + EOL" float2 aPixel = (float2) (theX / (float)theSizeX," + EOL" theY / (float)theSizeY);" + EOL + EOL" float4 aP0 = mix (theOrigins.lo.lo, theOrigins.lo.hi, aPixel.x);" + EOL" float4 aP1 = mix (theOrigins.hi.lo, theOrigins.hi.hi, aPixel.x);" + EOL + EOL" theRay->Origin = mix (aP0, aP1, aPixel.y);" + EOL + EOL" aP0 = mix (theDirects.lo.lo, theDirects.lo.hi, aPixel.x);" + EOL" aP1 = mix (theDirects.hi.lo, theDirects.hi.hi, aPixel.x);" + EOL + EOL" theRay->Direct = mix (aP0, aP1, aPixel.y);" + EOL" }" + EOL + EOL + ///////////////////////////////////////////////////////////////////////////////////////// + // Functions for compute ray-object intersection + EOL + EOL" #define _OOEPS_ exp2( -80.0f )" + EOL + // ======================================================================= + // function : IntersectSphere + // purpose : Computes ray-sphere intersection + // ======================================================================= + EOL" bool IntersectSphere (const SRay* theRay, float theRadius, float* theTime)" + EOL" {" + EOL" float aDdotD = dot (theRay->Direct.xyz, theRay->Direct.xyz);" + EOL" float aDdotO = dot (theRay->Direct.xyz, theRay->Origin.xyz);" + EOL" float aOdotO = dot (theRay->Origin.xyz, theRay->Origin.xyz);" + EOL + EOL" float aD = aDdotO * aDdotO - aDdotD * (aOdotO - theRadius * theRadius);" + EOL + EOL" if (aD > 0.f)" + EOL" {" + EOL" *theTime = (-aDdotO + native_sqrt (aD)) * (1.f / aDdotD);" + EOL + EOL" return *theTime > 0.f;" + EOL" }" + EOL + EOL" return false;" + EOL" }" + EOL + // ======================================================================= + // function : IntersectBox + // purpose : Computes ray-box intersection (slab test) + // ======================================================================= + EOL" bool IntersectBox (const SRay* theRay," + EOL" float4 theMinPoint," + EOL" float4 theMaxPoint," + EOL" float* theTimeStart," + EOL" float* theTimeFinal)" + EOL" {" + EOL" const float4 aInvDirect = (float4)(" + EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?" + EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x))," + EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?" + EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y))," + EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?" + EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z))," + EOL" 0.f);" + EOL + EOL" const float4 aTime0 = (theMinPoint - theRay->Origin) * aInvDirect;" + EOL" const float4 aTime1 = (theMaxPoint - theRay->Origin) * aInvDirect;" + EOL + EOL" const float4 aTimeMax = max (aTime0, aTime1);" + EOL" const float4 aTimeMin = min (aTime0, aTime1);" + EOL + EOL" *theTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" + EOL" *theTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" + EOL + EOL" return (*theTimeStart <= *theTimeFinal) & (*theTimeFinal >= 0.f);" + EOL" }" + EOL + // ======================================================================= + // function : IntersectNodes + // purpose : Computes intersection of ray with two child nodes (boxes) + // ======================================================================= + EOL" void IntersectNodes (const SRay* theRay," + EOL" float4 theMinPoint0," + EOL" float4 theMaxPoint0," + EOL" float4 theMinPoint1," + EOL" float4 theMaxPoint1," + EOL" float* theTimeStart0," + EOL" float* theTimeStart1," + EOL" float theMaxTime)" + EOL" {" + EOL" const float4 aInvDirect = (float4)(" + EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?" + EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x))," + EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?" + EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y))," + EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?" + EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z))," + EOL" 0.f);" + EOL + EOL" float4 aTime0 = (theMinPoint0 - theRay->Origin) * aInvDirect;" + EOL" float4 aTime1 = (theMaxPoint0 - theRay->Origin) * aInvDirect;" + EOL + EOL" float4 aTimeMax = max (aTime0, aTime1);" + EOL" float4 aTimeMin = min (aTime0, aTime1);" + EOL + EOL" aTime0 = (theMinPoint1 - theRay->Origin) * aInvDirect;" + EOL" aTime1 = (theMaxPoint1 - theRay->Origin) * aInvDirect;" + EOL + EOL" float aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" + EOL" float aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" + EOL + EOL" aTimeMax = max (aTime0, aTime1);" + EOL" aTimeMin = min (aTime0, aTime1);" + EOL + EOL" *theTimeStart0 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)" + EOL" ? aTimeStart : -MAXFLOAT;" + EOL + EOL" aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" + EOL" aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" + EOL + EOL" *theTimeStart1 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)" + EOL" ? aTimeStart : -MAXFLOAT;" + EOL" }" + EOL + // ======================================================================= + // function : IntersectTriangle + // purpose : Computes ray-triangle intersection (branchless version) + // ======================================================================= + EOL" bool IntersectTriangle (const SRay* theRay," + EOL" const float4 thePoint0," + EOL" const float4 thePoint1," + EOL" const float4 thePoint2," + EOL" float4* theNormal," + EOL" float* theTime," + EOL" float* theU," + EOL" float* theV)" + EOL" {" + EOL" const float4 aEdge0 = thePoint1 - thePoint0;" + EOL" const float4 aEdge1 = thePoint0 - thePoint2;" + EOL + EOL" *theNormal = cross (aEdge1, aEdge0);" + EOL + EOL" const float4 aEdge2 = (1.f / dot (*theNormal, theRay->Direct)) * (thePoint0 - theRay->Origin);" + EOL + EOL" *theTime = dot (*theNormal, aEdge2);" + EOL + EOL" const float4 theInc = cross (theRay->Direct, aEdge2);" + EOL + EOL" *theU = dot (theInc, aEdge1);" + EOL" *theV = dot (theInc, aEdge0);" + EOL + EOL" return (*theTime > 0) & (*theU >= 0.f) & (*theV >= 0.f) & (*theU + *theV <= 1.f);" + EOL" }" + EOL + ///////////////////////////////////////////////////////////////////////////////////////// + // Support shading functions + EOL + EOL" const sampler_t EnvironmentSampler =" + EOL" CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_REPEAT | CLK_FILTER_LINEAR;" + EOL + // ======================================================================= + // function : SmoothNormal + // purpose : Interpolates normal across the triangle + // ======================================================================= + EOL" float4 SmoothNormal (__global float4* theNormals," + EOL" const SIntersect* theHit," + EOL" const int4 theIndices)" + EOL" {" + EOL" float4 aNormal0 = theNormals[theIndices.x]," + EOL" aNormal1 = theNormals[theIndices.y]," + EOL" aNormal2 = theNormals[theIndices.z];" + EOL + EOL" return fast_normalize (aNormal1 * theHit->U +" + EOL" aNormal2 * theHit->V +" + EOL" aNormal0 * (1.f - theHit->U - theHit->V));" + EOL" }" + EOL + // ======================================================================= + // function : Shade + // purpose : Computes Phong-based illumination + // ======================================================================= + EOL" float4 Shade (__global float4* theMaterials," + EOL" const float4 theLight," + EOL" const float4 theView," + EOL" const float4 theNormal," + EOL" const float4 theIntens," + EOL" const float theTranspr," + EOL" const int theMatIndex)" + EOL" {" + EOL" float aLambert = dot (theNormal, theLight);" + EOL + EOL" aLambert = theTranspr > 0.f ? fabs (aLambert) : aLambert;" + EOL + EOL" if (aLambert > 0.f)" + EOL" {" + EOL" const float4 aMatDiff = theMaterials[7 * theMatIndex + 1];" + EOL" const float4 aMatSpec = theMaterials[7 * theMatIndex + 2];" + EOL + EOL" const float4 aReflect = 2.f * dot (theLight, theNormal) * theNormal - theLight;" + EOL + EOL" const float aSpecular = pow (max (dot (aReflect.xyz, theView.xyz), 0.f), aMatSpec.w);" + EOL + EOL" return theIntens * (aMatDiff * aLambert + aMatSpec * aSpecular);" + EOL" }" + EOL + EOL" return ZERO;" + EOL" }" + EOL + // ======================================================================= + // function : Lat-long + // purpose : Converts world direction to environment texture coordinates + // ======================================================================= + EOL" float2 Latlong (const float4 theDirect)" + EOL" {" + EOL" float aPsi = acos( -theDirect.y );" + EOL" float aPhi = atan2( theDirect.z, theDirect.x );" + EOL + EOL" aPhi = (aPhi < 0) ? (aPhi + 2.f * M_PI_F) : aPhi;" + EOL + EOL" return (float2) (aPhi / (2.f * M_PI_F), aPsi / M_PI_F);" + EOL" }" + EOL + ///////////////////////////////////////////////////////////////////////////////////////// + // Core ray tracing function + EOL + // ======================================================================= + // function : push + // purpose : Pushes BVH node index to local stack + // ======================================================================= + EOL" void push (uint* theStack, char* thePos, const uint theValue)" + EOL" {" + EOL" (*thePos)++;" + EOL" theStack[*thePos] = theValue;" + EOL" }" + EOL + // ======================================================================= + // function : pop + // purpose : Pops BVH node index from local stack + // ======================================================================= + EOL" void pop (uint* theStack, char* thePos, uint* theValue)" + EOL" {" + EOL" *theValue = theStack[*thePos];" + EOL" (*thePos)--;" + EOL" }" + EOL + // #define BVH_MINIMIZE_MEM_LOADS + EOL + // ======================================================================= + // function : Traverse + // purpose : Finds intersection with nearest triangle + // ======================================================================= + EOL" int4 Traverse (const SRay* theRay," + EOL" __global int4* theIndices," + EOL" __global float4* theVertices," + EOL" __global float4* theNodeMinPoints," + EOL" __global float4* theNodeMaxPoints," + EOL" __global int4* theNodeDataRecords," + EOL" SIntersect* theHit)" + EOL" {" + EOL" uint aStack [32];" + EOL" char aHead = -1;" + EOL + EOL" uint aNode = 0;" // root node + EOL + EOL" float aTimeMin1;" + EOL" float aTimeMin2;" + EOL + EOL" float4 aNodeMinLft;" + EOL" float4 aNodeMaxLft;" + EOL" float4 aNodeMinRgh;" + EOL" float4 aNodeMaxRgh;" + EOL + EOL" #ifdef BVH_MINIMIZE_MEM_LOADS" + EOL" aNodeMinLft = theNodeMinPoints[aNode];" + EOL" aNodeMaxLft = theNodeMaxPoints[aNode];" + EOL" #endif" + EOL + EOL" int4 aTriangleIndex = (int4) (-1);" + EOL + EOL" theHit->Time = MAXFLOAT;" + EOL + EOL" #ifdef BVH_MINIMIZE_MEM_LOADS" + EOL" int3 aData = (int3) (1," + EOL" as_int (aNodeMinLft.w)," + EOL" as_int (aNodeMaxLft.w));" + EOL + EOL" aData = aData.y < 0 ? -aData : aData;" + EOL" #endif" + EOL + EOL" while (true)" + EOL" {" + EOL" #ifndef BVH_MINIMIZE_MEM_LOADS" + EOL" int3 aData = theNodeDataRecords[aNode].xyz;" + EOL" #endif" + EOL + EOL" if (aData.x != 1)" // if inner node + EOL" {" + EOL" aNodeMinLft = theNodeMinPoints[aData.y];" + EOL" aNodeMinRgh = theNodeMinPoints[aData.z];" + EOL" aNodeMaxLft = theNodeMaxPoints[aData.y];" + EOL" aNodeMaxRgh = theNodeMaxPoints[aData.z];" + EOL + EOL" IntersectNodes (theRay," + EOL" aNodeMinLft," + EOL" aNodeMaxLft," + EOL" aNodeMinRgh," + EOL" aNodeMaxRgh," + EOL" &aTimeMin1," + EOL" &aTimeMin2," + EOL" theHit->Time);" + EOL + EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);" + EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);" + EOL + EOL" if (aHitLft & aHitRgh)" + EOL" {" + EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;" + EOL + EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);" + EOL + EOL" #ifdef BVH_MINIMIZE_MEM_LOADS" + EOL" aData = (int3) (1," + EOL" as_int (aTimeMin1 < aTimeMin2 ? aNodeMinLft.w : aNodeMinRgh.w)," + EOL" as_int (aTimeMin1 < aTimeMin2 ? aNodeMaxLft.w : aNodeMaxRgh.w));" + EOL + EOL" aData = aData.y < 0 ? -aData : aData;" + EOL" #endif" + EOL" }" + EOL" else" + EOL" {" + EOL" if (aHitLft | aHitRgh)" + EOL" {" + EOL" aNode = aHitLft ? aData.y : aData.z;" + EOL + EOL" #ifdef BVH_MINIMIZE_MEM_LOADS" + EOL" aData = (int3) (1," + EOL" as_int (aHitLft ? aNodeMinLft.w : aNodeMinRgh.w)," + EOL" as_int (aHitLft ? aNodeMaxLft.w : aNodeMaxRgh.w));" + EOL + EOL" aData = aData.y < 0 ? -aData : aData;" + EOL" #endif" + EOL" }" + EOL" else" + EOL" {" + EOL" if (aHead < 0)" + EOL" return aTriangleIndex;" + EOL + EOL" pop (aStack, &aHead, &aNode);" + EOL + EOL" #ifdef BVH_MINIMIZE_MEM_LOADS" + EOL" aData = theNodeDataRecords[aNode].xyz;" + EOL" #endif" + EOL" }" + EOL" }" + EOL" }" + EOL" else " // if leaf node + EOL" {" + EOL" for (int nTri = aData.y; nTri <= aData.z; ++nTri)" + EOL" {" + EOL" int4 anIndex = theIndices[nTri];" + EOL + EOL" const float4 aP0 = theVertices[anIndex.x];" + EOL" const float4 aP1 = theVertices[anIndex.y];" + EOL" const float4 aP2 = theVertices[anIndex.z];" + EOL + EOL" float4 aNormal;" + EOL + EOL" float aTime, aU, aV;" + EOL + EOL" if (IntersectTriangle (theRay, aP0, aP1, aP2, &aNormal, &aTime, &aU, &aV) & (aTime < theHit->Time))" + EOL" {" + EOL" aTriangleIndex = anIndex;" + EOL" theHit->Normal = aNormal;" + EOL" theHit->Time = aTime;" + EOL" theHit->U = aU;" + EOL" theHit->V = aV;" + EOL" }" + EOL" }" + EOL + EOL" if (aHead < 0)" + EOL" return aTriangleIndex;" + EOL + EOL" pop (aStack, &aHead, &aNode);" + EOL + EOL" #ifdef BVH_MINIMIZE_MEM_LOADS" + EOL" aData = theNodeDataRecords[aNode].xyz;" + EOL" #endif" + EOL" }" + EOL" }" + EOL + EOL" return aTriangleIndex;" + EOL" }" + EOL + EOL" #define TRANSPARENT_SHADOW_" + EOL + // ======================================================================= + // function : TraverseShadow + // purpose : Finds intersection with any triangle + // ======================================================================= + EOL" float TraverseShadow (const SRay* theRay," + EOL" __global int4* theIndices," + EOL" __global float4* theVertices," + EOL" __global float4* materials," + EOL" __global float4* theNodeMinPoints," + EOL" __global float4* theNodeMaxPoints," + EOL" __global int4* theNodeDataRecords," + EOL" float theDistance)" + EOL" {" + EOL" uint aStack [32];" + EOL" char aHead = -1;" + EOL + EOL" uint aNode = 0;" // root node + EOL + EOL" float aFactor = 1.f;" // light attenuation factor + EOL + EOL" float aTimeMin1;" + EOL" float aTimeMin2;" + EOL + EOL" while (true)" + EOL" {" + EOL" int3 aData = theNodeDataRecords[aNode].xyz;" + EOL + EOL" if (aData.x != 1)" // if inner node + EOL" {" + EOL" IntersectNodes (theRay," + EOL" theNodeMinPoints[aData.y]," + EOL" theNodeMaxPoints[aData.y]," + EOL" theNodeMinPoints[aData.z]," + EOL" theNodeMaxPoints[aData.z]," + EOL" &aTimeMin1," + EOL" &aTimeMin2," + EOL" theDistance);" + EOL + EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);" + EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);" + EOL + EOL" if (aHitLft & aHitRgh)" + EOL" {" + EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;" + EOL + EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);" + EOL" }" + EOL" else" + EOL" {" + EOL" if (aHitLft | aHitRgh)" + EOL" {" + EOL" aNode = aHitLft ? aData.y : aData.z;" + EOL" }" + EOL" else" + EOL" {" + EOL" if (aHead < 0)" + EOL" return aFactor;" + EOL + EOL" pop (aStack, &aHead, &aNode);" + EOL" }" + EOL" }" + EOL" }" + EOL" else " // if leaf node + EOL" {" + EOL" for (int nTri = aData.y; nTri <= aData.z; ++nTri)" + EOL" {" + EOL" int4 anIndex = theIndices[nTri];" + EOL + EOL" const float4 aP0 = theVertices[anIndex.x];" + EOL" const float4 aP1 = theVertices[anIndex.y];" + EOL" const float4 aP2 = theVertices[anIndex.z];" + EOL + EOL" float4 aNormal;" + EOL + EOL" float aTime, aU, aV;" + EOL + EOL" if (IntersectTriangle (theRay, aP0, aP1, aP2, &aNormal, &aTime, &aU, &aV) & (aTime < theDistance))" + EOL" {" + EOL" #ifdef TRANSPARENT_SHADOW" + EOL" aFactor *= materials[7 * index.w + 6].x;" + EOL + EOL" if (aFactor < 0.1f)" + EOL" return aFactor;" + EOL" #else" + EOL" return 0.f;" + EOL" #endif" + EOL" }" + EOL" }" + EOL + EOL" if (aHead < 0)" + EOL" return aFactor;" + EOL + EOL" pop (aStack, &aHead, &aNode);" + EOL" }" + EOL" }" + EOL + EOL" return aFactor;" + EOL" }" + EOL + EOL" #define _MAX_DEPTH_ 5" + EOL + EOL" #define _MAT_SIZE_ 7" + EOL + EOL" #define _LGH_SIZE_ 3" + EOL + // ======================================================================= + // function : Raytrace + // purpose : Computes color of specified ray + // ======================================================================= + EOL" float4 Raytrace (SRay* theRay," + EOL" __read_only image2d_t theEnvMap," + EOL" __global float4* theNodeMinPoints," + EOL" __global float4* theNodeMaxPoints," + EOL" __global int4* theNodeDataRecords," + EOL" __global float4* theLightSources," + EOL" __global float4* theMaterials," + EOL" __global float4* theVertices," + EOL" __global float4* theNormals," + EOL" __global int4* theIndices," + EOL" const int theLightCount," + EOL" const float theEpsilon," + EOL" const float theRadius," + EOL" const int isShadows," + EOL" const int isReflect)" + EOL" {" + EOL" float4 aResult = (float4) (0.f, 0.f, 0.f, 0.f);" + EOL" float4 aWeight = (float4) (1.f, 1.f, 1.f, 1.f);" + EOL + EOL" SIntersect aHit;" + EOL + EOL" for (int aDepth = 0; aDepth < _MAX_DEPTH_; ++aDepth)" + EOL" {" + EOL" int4 aTriangle = Traverse (theRay," + EOL" theIndices," + EOL" theVertices," + EOL" theNodeMinPoints," + EOL" theNodeMaxPoints," + EOL" theNodeDataRecords," + EOL" &aHit);" + EOL + EOL" if (aTriangle.x < 0.f)" + EOL" {" + EOL" float aTime;" + EOL + EOL" if (aWeight.w != 0.f || !IntersectSphere (theRay, theRadius, &aTime))" + EOL" break;" + EOL + EOL" float2 aTexCoord = Latlong (fma (theRay->Direct, (float4) (aTime), theRay->Origin) * (1.f / theRadius));" + EOL + EOL" aResult += aWeight * read_imagef (theEnvMap, EnvironmentSampler, aTexCoord);" + EOL + EOL" return (float4) (aResult.x," + EOL" aResult.y," + EOL" aResult.z," + EOL" aWeight.w);" + EOL" }" + EOL + EOL" " // Compute geometric normal + EOL" float4 aGeomNormal = aHit.Normal; aGeomNormal = fast_normalize (aGeomNormal);" + EOL + EOL" " // Compute interpolated normal + EOL" float4 aNormal = SmoothNormal (theNormals, &aHit, aTriangle);" + EOL + EOL" " // Compute intersection point + EOL" float4 aPoint = theRay->Direct * aHit.Time + theRay->Origin;" + EOL + EOL" float4 aMaterAmb = theMaterials [_MAT_SIZE_ * aTriangle.w + 0];" + EOL" float4 aMaterTrn = theMaterials [_MAT_SIZE_ * aTriangle.w + 6];" + EOL + EOL" for (int nLight = 0; nLight < theLightCount; ++nLight)" + EOL" {" + EOL" float4 aLightAmbient = theLightSources [_LGH_SIZE_ * nLight];" + EOL + EOL" aResult += aWeight * aLightAmbient * aMaterAmb *" + EOL" (aMaterTrn.x * max (fabs (dot (theRay->Direct, aNormal)), 0.5f));" + EOL + EOL" if (aLightAmbient.w < 0.f)" // 'ambient' light + EOL" {" + EOL" continue;" // 'ambient' light has no another luminances + EOL" }" + EOL + EOL" float4 aLightPosition = theLightSources [_LGH_SIZE_ * nLight + 2];" + EOL + EOL" SRay aShadow;" + EOL" aShadow.Direct = aLightPosition;" + EOL + EOL" float aLightDistance = MAXFLOAT;" + EOL" if (aLightPosition.w != 0.f)" + EOL" {" + EOL" aLightDistance = length (aLightPosition - aPoint);" + EOL" aShadow.Direct = (aLightPosition - aPoint) * (1.f / aLightDistance);" + EOL" }" + EOL + EOL" aShadow.Origin = aPoint + aShadow.Direct * theEpsilon +" + EOL" aGeomNormal * copysign (theEpsilon, dot (aGeomNormal, aShadow.Direct));" + EOL + EOL" float aFactor = 1.f;" + EOL + EOL" if (isShadows)" + EOL" {" + EOL" aFactor = TraverseShadow (&aShadow," + EOL" theIndices," + EOL" theVertices," + EOL" theMaterials," + EOL" theNodeMinPoints," + EOL" theNodeMaxPoints," + EOL" theNodeDataRecords," + EOL" aLightDistance);" + EOL" }" + EOL + EOL" aResult += (aMaterTrn.x * aFactor) * aWeight * Shade (theMaterials," + EOL" aShadow.Direct," + EOL" -theRay->Direct," + EOL" aNormal," + EOL" UNIT," + EOL" aMaterTrn.y," + EOL" aTriangle.w);" + EOL" }" + EOL + EOL" if (aMaterTrn.y > 0.f)" + EOL" {" + EOL" aWeight *= aMaterTrn.y;" + EOL" }" + EOL" else" + EOL" {" + EOL" float4 aMaterRef = theMaterials [_MAT_SIZE_ * aTriangle.w + 4];" + EOL" aWeight *= isReflect ? aMaterRef : ZERO;" + EOL + EOL" theRay->Direct -= 2.f * dot (theRay->Direct, aNormal) * aNormal;" + EOL + EOL" float aDdotN = dot (theRay->Direct, aGeomNormal);" + EOL" if (aDdotN < 0.f)" + EOL" theRay->Direct -= aDdotN * aGeomNormal;" + EOL" }" + EOL + EOL" if (aWeight.x < 0.1f && aWeight.y < 0.1f && aWeight.z < 0.1f)" + EOL" {" + EOL" return (float4) (aResult.x," + EOL" aResult.y," + EOL" aResult.z," + EOL" aWeight.w);" + EOL" }" + EOL + EOL" theRay->Origin = theRay->Direct * theEpsilon + aPoint;" + EOL" }" + EOL + EOL" return (float4) (aResult.x," + EOL" aResult.y," + EOL" aResult.z," + EOL" aWeight.w);" + EOL" }" + EOL + EOL + /////////////////////////////////////////////////////////////////////////////// + // Ray tracing kernel functions + EOL + // ======================================================================= + // function : Main + // purpose : Computes pixel color using ray-tracing + // ======================================================================= + EOL" __kernel void Main (__write_only image2d_t theOutput," + EOL" __read_only image2d_t theEnvMap," + EOL" __global float4* theNodeMinPoints," + EOL" __global float4* theNodeMaxPoints," + EOL" __global int4* theNodeDataRecords," + EOL" __global float4* theLightSources," + EOL" __global float4* theMaterials," + EOL" __global float4* theVertices," + EOL" __global float4* theNormals," + EOL" __global int4* theIndices," + EOL" const float16 theOrigins," + EOL" const float16 theDirects," + EOL" const int theLightCount," + EOL" const float theEpsilon," + EOL" const float theRadius," + EOL" const int isShadows," + EOL" const int isReflect," + EOL" const int theSizeX," + EOL" const int theSizeY)" + EOL" {" + EOL" const int aX = get_global_id (0);" + EOL" const int aY = get_global_id (1);" + EOL + EOL" if (aX >= theSizeX || aY >= theSizeY)" + EOL" return;" + EOL + EOL" private SRay aRay;" + EOL + EOL" GenerateRay (&aRay," + EOL" aX," + EOL" aY," + EOL" theSizeX," + EOL" theSizeY," + EOL" theOrigins," + EOL" theDirects);" + EOL + EOL" float4 aBoxMin = theNodeMinPoints[0] - (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);" + EOL" float4 aBoxMax = theNodeMaxPoints[0] + (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);" + EOL + EOL" float aTimeStart;" + EOL" float aTimeFinal;" + EOL + EOL" float4 aColor = (float4) (0.f, 0.f, 0.f, 1.f);" + EOL + EOL" if (IntersectBox (&aRay, aBoxMin, aBoxMax, &aTimeStart, &aTimeFinal))" + EOL" {" + EOL" aRay.Origin = fma (aRay.Direct, (float4) (aTimeStart), aRay.Origin);" + EOL + EOL" aColor = Raytrace (&aRay," + EOL" theEnvMap," + EOL" theNodeMinPoints," + EOL" theNodeMaxPoints," + EOL" theNodeDataRecords," + EOL" theLightSources," + EOL" theMaterials," + EOL" theVertices," + EOL" theNormals," + EOL" theIndices," + EOL" theLightCount," + EOL" theEpsilon," + EOL" theRadius," + EOL" isShadows," + EOL" isReflect);" + EOL" }" + EOL + EOL" write_imagef (theOutput, (int2) (aX, aY), aColor);" + EOL" }" + EOL + EOL" const sampler_t OutputSampler =" + EOL" CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;" + EOL + EOL" #define _LUM_DELTA_ 0.075f" + EOL + EOL" #define AA_MAX 0.559017f" + EOL" #define AA_MIN 0.186339f" + EOL + // ======================================================================= + // function : MainAntialiased + // purpose : Performs adaptive sub-pixel rendering + // ======================================================================= + EOL" __kernel void MainAntialiased ( __read_only image2d_t theInput," + EOL" __write_only image2d_t theOutput," + EOL" __read_only image2d_t theEnvMap," + EOL" __global float4* theNodeMinPoints," + EOL" __global float4* theNodeMaxPoints," + EOL" __global int4* theNodeDataRecords," + EOL" __global float4* theLightSources," + EOL" __global float4* theMaterials," + EOL" __global float4* theVertices," + EOL" __global float4* theNormals," + EOL" __global int4* theIndices," + EOL" const float16 theOrigins," + EOL" const float16 theDirects," + EOL" const int theLightCount," + EOL" const float theEpsilon," + EOL" const float theRadius," + EOL" const int isShadows," + EOL" const int isReflect," + EOL" const int theSizeX," + EOL" const int theSizeY )" + EOL" {" + EOL" const int aX = get_global_id (0);" + EOL" const int aY = get_global_id (1);" + EOL + EOL" if (aX >= theSizeX || aY >= theSizeY)" + EOL" return;" + EOL + EOL" float4 aClr0 = read_imagef (theInput, OutputSampler, (float2) (aX + 0, aY + 0));" + EOL" float4 aClr1 = read_imagef (theInput, OutputSampler, (float2) (aX + 0, aY - 1));" + EOL" float4 aClr2 = read_imagef (theInput, OutputSampler, (float2) (aX + 0, aY + 1));" + EOL + EOL" float4 aClr3 = read_imagef (theInput, OutputSampler, (float2) (aX + 1, aY + 0));" + EOL" float4 aClr4 = read_imagef (theInput, OutputSampler, (float2) (aX + 1, aY - 1));" + EOL" float4 aClr5 = read_imagef (theInput, OutputSampler, (float2) (aX + 1, aY + 1));" + EOL + EOL" float4 aClr6 = read_imagef (theInput, OutputSampler, (float2) (aX - 1, aY + 0));" + EOL" float4 aClr7 = read_imagef (theInput, OutputSampler, (float2) (aX - 1, aY - 1));" + EOL" float4 aClr8 = read_imagef (theInput, OutputSampler, (float2) (aX - 1, aY + 1));" + EOL + EOL" aClr1 = (aClr1.w == 1.f) ? -UNIT : aClr1;" + EOL" aClr2 = (aClr2.w == 1.f) ? -UNIT : aClr2;" + EOL" aClr3 = (aClr3.w == 1.f) ? -UNIT : aClr3;" + EOL" aClr4 = (aClr4.w == 1.f) ? -UNIT : aClr4;" + EOL" aClr5 = (aClr5.w == 1.f) ? -UNIT : aClr5;" + EOL" aClr6 = (aClr6.w == 1.f) ? -UNIT : aClr6;" + EOL" aClr7 = (aClr7.w == 1.f) ? -UNIT : aClr7;" + EOL" aClr8 = (aClr8.w == 1.f) ? -UNIT : aClr8;" + EOL + EOL" float aLum = (aClr0.w == 1.f) ? -1.f : (0.2126f * aClr0.x + 0.7152f * aClr0.y + 0.0722f * aClr0.z);" + EOL + EOL + EOL" bool render = fabs (0.2126f * aClr1.x + 0.7152f * aClr1.y + 0.0722f * aClr1.z - aLum) > _LUM_DELTA_ ||" + EOL" fabs (0.2126f * aClr2.x + 0.7152f * aClr2.y + 0.0722f * aClr2.z - aLum) > _LUM_DELTA_ ||" + EOL" fabs (0.2126f * aClr3.x + 0.7152f * aClr3.y + 0.0722f * aClr3.z - aLum) > _LUM_DELTA_ ||" + EOL" fabs (0.2126f * aClr4.x + 0.7152f * aClr4.y + 0.0722f * aClr4.z - aLum) > _LUM_DELTA_ ||" + EOL" fabs (0.2126f * aClr5.x + 0.7152f * aClr5.y + 0.0722f * aClr5.z - aLum) > _LUM_DELTA_ ||" + EOL" fabs (0.2126f * aClr6.x + 0.7152f * aClr6.y + 0.0722f * aClr6.z - aLum) > _LUM_DELTA_ ||" + EOL" fabs (0.2126f * aClr7.x + 0.7152f * aClr7.y + 0.0722f * aClr7.z - aLum) > _LUM_DELTA_ ||" + EOL" fabs (0.2126f * aClr8.x + 0.7152f * aClr8.y + 0.0722f * aClr8.z - aLum) > _LUM_DELTA_;" + EOL + EOL" float4 aColor = aClr0;" + EOL + EOL" private SRay aRay;" + EOL + EOL" const float4 aBoxMin = theNodeMinPoints[0] - (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);" + EOL" const float4 aBoxMax = theNodeMaxPoints[0] + (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);" + EOL + EOL" if (render)" + EOL" {" + EOL" for (int aSample = 0; aSample <= 3; ++aSample)" + EOL" {" + EOL" float fX = aX, fY = aY;" + EOL + EOL" if (aSample == 0)" + EOL" {" + EOL" fX -= AA_MIN; fY -= AA_MAX;" + EOL" }" + EOL" else if (aSample == 1)" + EOL" {" + EOL" fX -= AA_MAX; fY += AA_MIN;" + EOL" }" + EOL" else if (aSample == 2)" + EOL" {" + EOL" fX += AA_MIN; fY += AA_MAX;" + EOL" }" + EOL" else" + EOL" {" + EOL" fX += AA_MAX; fY -= AA_MIN;" + EOL" }" + EOL + EOL" GenerateRay (&aRay," + EOL" fX," + EOL" fY," + EOL" theSizeX," + EOL" theSizeY," + EOL" theOrigins," + EOL" theDirects);" + EOL + EOL" float aTimeStart;" + EOL" float aTimeFinal;" + EOL + EOL" if (IntersectBox (&aRay, aBoxMin, aBoxMax, &aTimeStart, &aTimeFinal))" + EOL" {" + EOL" aRay.Origin = fma (aRay.Direct, (float4) (aTimeStart), aRay.Origin);" + EOL + EOL" aColor += Raytrace (&aRay," + EOL" theEnvMap," + EOL" theNodeMinPoints," + EOL" theNodeMaxPoints," + EOL" theNodeDataRecords," + EOL" theLightSources," + EOL" theMaterials," + EOL" theVertices," + EOL" theNormals," + EOL" theIndices," + EOL" theLightCount," + EOL" theEpsilon," + EOL" theRadius," + EOL" isShadows," + EOL" isReflect);" + EOL" }" + EOL" else" + EOL" aColor += (float4) (0.f, 0.f, 0.f, 1.f);" + EOL" }" + EOL + EOL" aColor *= 1.f / 5.f;" + EOL" }" + EOL + EOL" write_imagef (theOutput, (int2) (aX, aY), aColor);" + EOL" }"; + +#endif diff --git a/src/OpenGl/OpenGl_RaytraceTypes.hxx b/src/OpenGl/OpenGl_RaytraceTypes.hxx new file mode 100644 index 0000000000..c75ae3f6c5 --- /dev/null +++ b/src/OpenGl/OpenGl_RaytraceTypes.hxx @@ -0,0 +1,46 @@ +// Created on: 2013-10-15 +// Created by: Denis BOGOLEPOV +// Copyright (c) 2013 OPEN CASCADE SAS +// +// The content of this file is subject to the Open CASCADE Technology Public +// License Version 6.5 (the "License"). You may not use the content of this file +// except in compliance with the License. Please obtain a copy of the License +// at http://www.opencascade.org and read it completely before using this file. +// +// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its +// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France. +// +// The Original Code and all software distributed under the License is +// distributed on an "AS IS" basis, without warranty of any kind, and the +// Initial Developer hereby disclaims all such warranties, including without +// limitation, any warranties of merchantability, fitness for a particular +// purpose or non-infringement. Please see the License for the specific terms +// and conditions governing the rights and limitations under the License. + + +#ifndef _OpenGl_RaytraceTypes_Header +#define _OpenGl_RaytraceTypes_Header + +#include + +#include +#include + +//! 4D vector of integers. +typedef NCollection_Vec4 OpenGl_RTVec4i; + +//! 4D vector of floats. +typedef NCollection_Vec4 OpenGl_RTVec4f; + +//! 4D vector of doubles. +typedef NCollection_Vec4 OpenGl_RTVec4d; + +//! Array of 4D integer vectors. +typedef std::vector > OpenGl_RTArray4i; + +//! Array of 4D floating point vectors. +typedef std::vector > OpenGl_RTArray4f; + +#endif diff --git a/src/OpenGl/OpenGl_SceneGeometry.cxx b/src/OpenGl/OpenGl_SceneGeometry.cxx new file mode 100644 index 0000000000..5184e8dba9 --- /dev/null +++ b/src/OpenGl/OpenGl_SceneGeometry.cxx @@ -0,0 +1,761 @@ +// Created on: 2013-08-27 +// Created by: Denis BOGOLEPOV +// Copyright (c) 2013 OPEN CASCADE SAS +// +// The content of this file is subject to the Open CASCADE Technology Public +// License Version 6.5 (the "License"). You may not use the content of this file +// except in compliance with the License. Please obtain a copy of the License +// at http://www.opencascade.org and read it completely before using this file. +// +// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its +// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France. +// +// The Original Code and all software distributed under the License is +// distributed on an "AS IS" basis, without warranty of any kind, and the +// Initial Developer hereby disclaims all such warranties, including without +// limitation, any warranties of merchantability, fitness for a particular +// purpose or non-infringement. Please see the License for the specific terms +// and conditions governing the rights and limitations under the License. + +#ifdef HAVE_CONFIG_H + #include +#endif + +#ifdef HAVE_OPENCL + +#include + +#include + +namespace +{ + + //! Number of node bins per axis. + static const int THE_NUMBER_OF_BINS = 32; + + //! Max number of triangles per leaf node. + static const int THE_MAX_LEAF_TRIANGLES = 4; + + //! Useful constant for null integer 4D vector. + static const OpenGl_RTVec4i THE_ZERO_VEC_4I; + + //! Useful constant for null floating-point 4D vector. + static const OpenGl_RTVec4f THE_ZERO_VEC_4F; + +}; + +// ======================================================================= +// function : OpenGl_Material +// purpose : Creates new default material +// ======================================================================= +OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial() +: Ambient (THE_ZERO_VEC_4F), + Diffuse (THE_ZERO_VEC_4F), + Specular (THE_ZERO_VEC_4F), + Emission (THE_ZERO_VEC_4F), + Reflection (THE_ZERO_VEC_4F), + Refraction (THE_ZERO_VEC_4F), + Transparency (THE_ZERO_VEC_4F) +{ } + +// ======================================================================= +// function : OpenGl_Material +// purpose : Creates new material with specified properties +// ======================================================================= +OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient, + const OpenGl_RTVec4f& theDiffuse, + const OpenGl_RTVec4f& theSpecular) +: Ambient (theAmbient), + Diffuse (theDiffuse), + Specular (theSpecular), + Emission (THE_ZERO_VEC_4F), + Reflection (THE_ZERO_VEC_4F), + Refraction (THE_ZERO_VEC_4F), + Transparency (THE_ZERO_VEC_4F) +{ + // +} + +// ======================================================================= +// function : OpenGl_Material +// purpose : Creates new material with specified properties +// ======================================================================= +OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient, + const OpenGl_RTVec4f& theDiffuse, + const OpenGl_RTVec4f& theSpecular, + const OpenGl_RTVec4f& theEmission, + const OpenGl_RTVec4f& theTranspar) +: Ambient (theAmbient), + Diffuse (theDiffuse), + Specular (theSpecular), + Emission (theEmission), + Reflection (THE_ZERO_VEC_4F), + Refraction (THE_ZERO_VEC_4F), + Transparency (theTranspar) +{ + // +} + +// ======================================================================= +// function : OpenGl_Material +// purpose : Creates new material with specified properties +// ======================================================================= +OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient, + const OpenGl_RTVec4f& theDiffuse, + const OpenGl_RTVec4f& theSpecular, + const OpenGl_RTVec4f& theEmission, + const OpenGl_RTVec4f& theTranspar, + const OpenGl_RTVec4f& theReflection, + const OpenGl_RTVec4f& theRefraction) +: Ambient (theAmbient), + Diffuse (theDiffuse), + Specular (theSpecular), + Emission (theEmission), + Reflection (theReflection), + Refraction (theRefraction), + Transparency (theTranspar) +{ + // +} + +// ======================================================================= +// function : OpenGl_LightSource +// purpose : Creates new light source +// ======================================================================= +OpenGl_RaytraceLight::OpenGl_RaytraceLight (const OpenGl_RTVec4f& theAmbient) +: Ambient (theAmbient) +{ + // +} + +// ======================================================================= +// function : OpenGl_LightSource +// purpose : Creates new light source +// ======================================================================= +OpenGl_RaytraceLight::OpenGl_RaytraceLight (const OpenGl_RTVec4f& theDiffuse, + const OpenGl_RTVec4f& thePosition) +: Diffuse (theDiffuse), + Position (thePosition) +{ + // +} + +// ======================================================================= +// function : Center +// purpose : Returns centroid of specified triangle +// ======================================================================= +OpenGl_RTVec4f OpenGl_RaytraceScene::Center (const int theTriangle) const +{ + const OpenGl_RTVec4i anIndex (Triangles [theTriangle]); + return ( Vertices[anIndex.x()] + + Vertices[anIndex.y()] + + Vertices[anIndex.z()] ) * ( 1.f / 3.f ); +} + +// ======================================================================= +// function : CenterAxis +// purpose : Returns centroid of specified triangle +// ======================================================================= +float OpenGl_RaytraceScene::CenterAxis (const int theTriangle, + const int theAxis) const +{ + const OpenGl_RTVec4i anIndex (Triangles [theTriangle]); + return ( Vertices[anIndex.x()][theAxis] + + Vertices[anIndex.y()][theAxis] + + Vertices[anIndex.z()][theAxis] ) * ( 1.f / 3.f ); +} + +// ======================================================================= +// function : Box +// purpose : Returns AABB of specified triangle +// ======================================================================= +OpenGl_AABB OpenGl_RaytraceScene::Box (const int theTriangle) const +{ + const OpenGl_RTVec4i anIndex (Triangles[theTriangle]); + + const OpenGl_RTVec4f pA = Vertices[anIndex.x()]; + const OpenGl_RTVec4f pB = Vertices[anIndex.y()]; + const OpenGl_RTVec4f pC = Vertices[anIndex.z()]; + + OpenGl_AABB aBox (pA); + + aBox.Add (pB); + aBox.Add (pC); + + return aBox; +} + +// ======================================================================= +// function : Clear +// purpose : Clears all scene geometry data +// ======================================================================= +void OpenGl_RaytraceScene::Clear() +{ + AABB.Clear(); + + OpenGl_RTArray4f anEmptyNormals; + Normals.swap (anEmptyNormals); + + OpenGl_RTArray4f anEmptyVertices; + Vertices.swap (anEmptyVertices); + + OpenGl_RTArray4i anEmptyTriangles; + Triangles.swap (anEmptyTriangles); + + std::vector > anEmptyMaterials; + + Materials.swap (anEmptyMaterials); +} + +// ======================================================================= +// function : OpenGl_Node +// purpose : Creates new empty BVH node +// ======================================================================= +OpenGl_BVHNode::OpenGl_BVHNode() +: myMinPoint (THE_ZERO_VEC_4F), + myMaxPoint (THE_ZERO_VEC_4F), + myDataRcrd (THE_ZERO_VEC_4I) +{ + // +} + +// ======================================================================= +// function : OpenGl_Node +// purpose : Creates new BVH node with specified data +// ======================================================================= +OpenGl_BVHNode::OpenGl_BVHNode (const OpenGl_RTVec4f& theMinPoint, + const OpenGl_RTVec4f& theMaxPoint, + const OpenGl_RTVec4i& theDataRcrd) +: myMinPoint (theMinPoint), + myMaxPoint (theMaxPoint), + myDataRcrd (theDataRcrd) +{ + // +} + +// ======================================================================= +// function : OpenGl_Node +// purpose : Creates new leaf BVH node with specified data +// ======================================================================= +OpenGl_BVHNode::OpenGl_BVHNode (const OpenGl_AABB& theAABB, + const int theBegTriangle, + const int theEndTriangle) +: myMinPoint (theAABB.CornerMin()), + myMaxPoint (theAABB.CornerMax()), + myDataRcrd (1, + theBegTriangle, + theEndTriangle, + 0) +{ + // +} + +// ======================================================================= +// function : OpenGl_Node +// purpose : Creates new leaf BVH node with specified data +// ======================================================================= +OpenGl_BVHNode::OpenGl_BVHNode (const OpenGl_RTVec4f& theMinPoint, + const OpenGl_RTVec4f& theMaxPoint, + const int theBegTriangle, + const int theEndTriangle) +: myMinPoint (theMinPoint), + myMaxPoint (theMaxPoint), + myDataRcrd (1, + theBegTriangle, + theEndTriangle, + 0) +{ + // +} + +// ======================================================================= +// function : CleanUp +// purpose : Removes all tree nodes +// ======================================================================= +void OpenGl_BVH::CleanUp() +{ + OpenGl_RTArray4f anEmptyMinPointBuffer; + myMinPointBuffer.swap (anEmptyMinPointBuffer); + + OpenGl_RTArray4f anEmptyMaxPointBuffer; + myMaxPointBuffer.swap (anEmptyMaxPointBuffer); + + OpenGl_RTArray4i anEmptyDataRcrdBuffer; + myDataRcrdBuffer.swap (anEmptyDataRcrdBuffer); +} + +// ======================================================================= +// function : Node +// purpose : Returns node with specified index +// ======================================================================= +OpenGl_BVHNode OpenGl_BVH::Node (const int theIndex) const +{ + return OpenGl_BVHNode (myMinPointBuffer[theIndex], + myMaxPointBuffer[theIndex], + myDataRcrdBuffer[theIndex]); +} + +// ======================================================================= +// function : SetNode +// purpose : Replaces node with specified index +// ======================================================================= +void OpenGl_BVH::SetNode (const int theIndex, + const OpenGl_BVHNode& theNode) +{ + if (theIndex < static_cast (myMinPointBuffer.size())) + { + myMinPointBuffer[theIndex] = theNode.myMinPoint; + myMaxPointBuffer[theIndex] = theNode.myMaxPoint; + myDataRcrdBuffer[theIndex] = theNode.myDataRcrd; + } +} + +// ======================================================================= +// function : PushNode +// purpose : Adds new node to the tree +// ======================================================================= +int OpenGl_BVH::PushNode (const OpenGl_BVHNode& theNode) +{ + myMinPointBuffer.push_back (theNode.myMinPoint); + myMaxPointBuffer.push_back (theNode.myMaxPoint); + myDataRcrdBuffer.push_back (theNode.myDataRcrd); + return static_cast (myDataRcrdBuffer.size() - 1); +} + +// ======================================================================= +// function : OpenGl_NodeBuildTask +// purpose : Creates new node building task +// ======================================================================= +OpenGl_BVHNodeTask::OpenGl_BVHNodeTask() +: NodeToBuild (0), + BegTriangle (0), + EndTriangle (0) +{ + // +} + +// ======================================================================= +// function : OpenGl_NodeBuildTask +// purpose : Creates new node building task +// ======================================================================= +OpenGl_BVHNodeTask::OpenGl_BVHNodeTask (const int theNodeToBuild, + const int theBegTriangle, + const int theEndTriangle) +: NodeToBuild (theNodeToBuild), + BegTriangle (theBegTriangle), + EndTriangle (theEndTriangle) +{ + // +} + +// ======================================================================= +// function : OpenGl_BinnedBVHBuilder +// purpose : Creates new binned BVH builder +// ======================================================================= +OpenGl_BinnedBVHBuilder::OpenGl_BinnedBVHBuilder() +: myMaxDepth (30) +{ + // +} + +// ======================================================================= +// function : ~OpenGl_BinnedBVHBuilder +// purpose : Releases binned BVH builder +// ======================================================================= +OpenGl_BinnedBVHBuilder::~OpenGl_BinnedBVHBuilder() +{ + // +} + +#define BVH_DEBUG_OUTPUT_ + +#if defined( BVH_DEBUG_OUTPUT ) + #include +#endif + +// ======================================================================= +// function : ReinterpretIntAsFloat +// purpose : Reinterprets bits of integer value as floating-point value +// ======================================================================= +inline float ReinterpretIntAsFloat (int theValue) +{ + return *reinterpret_cast< float* > (&theValue); +} + +// ======================================================================= +// function : Build +// purpose : Builds BVH tree using binned SAH algorithm +// ======================================================================= +void OpenGl_BinnedBVHBuilder::Build (OpenGl_RaytraceScene& theGeometry, + const float theEpsilon) +{ + CleanUp(); + +#ifdef BVH_DEBUG_OUTPUT + std::cout << "Start building BVH..." << std::endl; + + std::cout << "Triangles: " << theGeometry.Triangles.size() << std::endl; +#endif + + if (theGeometry.Triangles.size() == 0) + return; + + // Create root node with all scene triangles + OpenGl_AABB anAABB = theGeometry.AABB; + anAABB.CornerMin() = OpenGl_RTVec4f (anAABB.CornerMin().x() - theEpsilon, + anAABB.CornerMin().y() - theEpsilon, + anAABB.CornerMin().z() - theEpsilon, + 1.0f); + anAABB.CornerMax() = OpenGl_RTVec4f (anAABB.CornerMax().x() + theEpsilon, + anAABB.CornerMax().y() + theEpsilon, + anAABB.CornerMax().z() + theEpsilon, + 1.0f); + myTree.PushNode (OpenGl_BVHNode (anAABB, 0, static_cast (theGeometry.Triangles.size() - 1))); + +#ifdef BVH_DEBUG_OUTPUT + std::cout << "Push root node: [" << 0 << ", " << + theGeometry.Triangles.size() - 1 << "]" << std::endl; +#endif + + // Setup splitting task for the root node + myNodeTasksQueue.push_back (OpenGl_BVHNodeTask (0, 0, static_cast (theGeometry.Triangles.size() - 1))); + + // Building nodes while tasks queue is not empty + for (int aTaskId = 0; aTaskId < static_cast (myNodeTasksQueue.size()); ++aTaskId) + { + BuildNode (theGeometry, aTaskId); + } + + // Write support data to optimize traverse + for (int aNode = 0; aNode < static_cast (myTree.DataRcrdBuffer().size()); ++aNode) + { + OpenGl_RTVec4i aData = myTree.DataRcrdBuffer()[aNode]; + myTree.MinPointBuffer()[aNode].w() = ReinterpretIntAsFloat (aData[0] ? aData[1] : -aData[1]); + myTree.MaxPointBuffer()[aNode].w() = ReinterpretIntAsFloat (aData[0] ? aData[2] : -aData[2]); + } +} + +// ======================================================================= +// function : CleanUp +// purpose : Clears previously built tree +// ======================================================================= +void OpenGl_BinnedBVHBuilder::CleanUp() +{ + myTree.CleanUp(); + myNodeTasksQueue.clear(); +} + +// ======================================================================= +// function : SetMaxDepth +// purpose : Sets maximum tree depth +// ======================================================================= +void OpenGl_BinnedBVHBuilder::SetMaxDepth (const int theMaxDepth) +{ + if (theMaxDepth > 1 && theMaxDepth < 30) + { + myMaxDepth = theMaxDepth - 1; + } +} + +//! Minimum node size to split. +static const float THE_NODE_MIN_SIZE = 1e-4f; + +// ======================================================================= +// function : BuildNode +// purpose : Builds node using task info +// ======================================================================= +void OpenGl_BinnedBVHBuilder::BuildNode (OpenGl_RaytraceScene& theGeometry, + const int theTask) +{ + OpenGl_BVHNodeTask aTask = myNodeTasksQueue[theTask]; + OpenGl_BVHNode aNode = myTree.Node (aTask.NodeToBuild); + +#ifdef BVH_DEBUG_OUTPUT + std::cout << "Build node " << aTask.NodeToBuild << ": [" << + aTask.BegTriangle << ", " << aTask.EndTriangle << "]" << std::endl; +#endif + + OpenGl_AABB anAABB (aNode.MinPoint(), aNode.MaxPoint()); + const OpenGl_RTVec4f aNodeSize = anAABB.Size(); + const float aNodeArea = anAABB.Area(); + + // Parameters for storing best split + float aMinSplitCost = std::numeric_limits::max(); + + int aMinSplitAxis = -1; + int aMinSplitIndex = 0; + int aMinSplitLftCount = 0; + int aMinSplitRghCount = 0; + + OpenGl_AABB aMinSplitLftAABB; + OpenGl_AABB aMinSplitRghAABB; + + // Find best split + for (int anAxis = 0; anAxis < 3; ++anAxis) + { + if (aNodeSize[anAxis] <= THE_NODE_MIN_SIZE) + continue; + + OpenGl_BinVector aBins (THE_NUMBER_OF_BINS); + GetSubVolumes (theGeometry, aNode, aBins, anAxis); + + // Choose the best split (with minimum SAH cost) + for (int aSplit = 1; aSplit < THE_NUMBER_OF_BINS; ++aSplit) + { + int aLftCount = 0; + int aRghCount = 0; + OpenGl_AABB aLftAABB; + OpenGl_AABB aRghAABB; + for (int anIndex = 0; anIndex < aSplit; ++anIndex) + { + aLftCount += aBins[anIndex].Count; + aLftAABB.Combine (aBins[anIndex].Volume); + } + + for (int anIndex = aSplit; anIndex < THE_NUMBER_OF_BINS; ++anIndex) + { + aRghCount += aBins[anIndex].Count; + aRghAABB.Combine (aBins[anIndex].Volume); + } + + // Simple SAH evaluation + float aCost = ( aLftAABB.Area() / aNodeArea ) * aLftCount + + ( aRghAABB.Area() / aNodeArea ) * aRghCount; + +#ifdef BVH_DEBUG_OUTPUT + std::cout << "\t\tBin " << aSplit << ", Cost = " << aCost << std::endl; +#endif + + if (aCost <= aMinSplitCost) + { + aMinSplitCost = aCost; + aMinSplitAxis = anAxis; + aMinSplitIndex = aSplit; + aMinSplitLftAABB = aLftAABB; + aMinSplitRghAABB = aRghAABB; + aMinSplitLftCount = aLftCount; + aMinSplitRghCount = aRghCount; + } + } + } + + if (aMinSplitAxis == -1) + { + // make outer (leaf) node + myTree.DataRcrdBuffer()[aTask.NodeToBuild].x() = 1; + return; + } + +#ifdef BVH_DEBUG_OUTPUT + switch (aMinSplitAxis) + { + case 0: + std::cout << "\tSplit axis: X = " << aMinSplitIndex << std::endl; + break; + case 1: + std::cout << "\tSplit axis: Y = " << aMinSplitIndex << std::endl; + break; + case 2: + std::cout << "\tSplit axis: Z = " << aMinSplitIndex << std::endl; + break; + } +#endif + + int aMiddle = SplitTriangles (theGeometry, aTask.BegTriangle, aTask.EndTriangle, + aNode, aMinSplitIndex - 1, aMinSplitAxis); + +#ifdef BVH_DEBUG_OUTPUT + std::cout << "\tLeft child: [" << aTask.BegTriangle << ", " + << aMiddle - 1 << "]" << std::endl; + + std::cout << "\tRight child: [" << aMiddle << ", " + << aTask.EndTriangle << "]" << std::endl; +#endif + +#define BVH_SIDE_LFT 1 +#define BVH_SIDE_RGH 2 + + // Setting up tasks for child nodes + for (int aSide = BVH_SIDE_LFT; aSide <= BVH_SIDE_RGH; ++aSide) + { + OpenGl_RTVec4f aMinPoint = (aSide == BVH_SIDE_LFT) + ? aMinSplitLftAABB.CornerMin() + : aMinSplitRghAABB.CornerMin(); + OpenGl_RTVec4f aMaxPoint = (aSide == BVH_SIDE_LFT) + ? aMinSplitLftAABB.CornerMax() + : aMinSplitRghAABB.CornerMax(); + + int aBegTriangle = (aSide == BVH_SIDE_LFT) + ? aTask.BegTriangle + : aMiddle; + int aEndTriangle = (aSide == BVH_SIDE_LFT) + ? aMiddle - 1 + : aTask.EndTriangle; + + OpenGl_BVHNode aChild (aMinPoint, aMaxPoint, aBegTriangle, aEndTriangle); + aChild.SetLevel (aNode.Level() + 1); + + // Check to see if child node must be split + const int aNbTriangles = (aSide == BVH_SIDE_LFT) + ? aMinSplitLftCount + : aMinSplitRghCount; + + const int isChildALeaf = (aNbTriangles <= THE_MAX_LEAF_TRIANGLES) || (aNode.Level() >= myMaxDepth); + if (isChildALeaf) + aChild.SetOuter(); + else + aChild.SetInner(); + + const int aChildIndex = myTree.PushNode (aChild); + + // Modify parent node + myTree.DataRcrdBuffer()[aTask.NodeToBuild].x() = 0; // inner node flag + if (aSide == BVH_SIDE_LFT) + myTree.DataRcrdBuffer()[aTask.NodeToBuild].y() = aChildIndex; // left child + else + myTree.DataRcrdBuffer()[aTask.NodeToBuild].z() = aChildIndex; // right child + + // Make new building task + if (!isChildALeaf) + myNodeTasksQueue.push_back (OpenGl_BVHNodeTask (aChildIndex, aBegTriangle, aEndTriangle)); + } +} + +// ======================================================================= +// function : SplitTriangles +// purpose : Splits node triangles into two intervals for child nodes +// ======================================================================= +int OpenGl_BinnedBVHBuilder::SplitTriangles (OpenGl_RaytraceScene& theGeometry, + const int theBegTriangle, + const int theEndTriangle, + OpenGl_BVHNode& theNode, + int theBin, + const int theAxis) +{ + int aLftIndex (theBegTriangle); + int aRghIndex (theEndTriangle); + + const float aMin = theNode.MinPoint()[theAxis]; + const float aMax = theNode.MaxPoint()[theAxis]; + + const float aStep = (aMax - aMin) / THE_NUMBER_OF_BINS; + + do + { + while ((int )floorf ((theGeometry.CenterAxis (aLftIndex, theAxis) - aMin) / aStep) <= theBin + && aLftIndex < theEndTriangle) + { + ++aLftIndex; + } + while ((int )floorf ((theGeometry.CenterAxis (aRghIndex, theAxis) - aMin) / aStep) > theBin + && aRghIndex > theBegTriangle) + { + --aRghIndex; + } + + if (aLftIndex <= aRghIndex) + { + if (aLftIndex != aRghIndex) + { + OpenGl_RTVec4i aLftTrg = theGeometry.Triangles[aLftIndex]; + OpenGl_RTVec4i aRghTrg = theGeometry.Triangles[aRghIndex]; + theGeometry.Triangles[aLftIndex] = aRghTrg; + theGeometry.Triangles[aRghIndex] = aLftTrg; + } + + aLftIndex++; aRghIndex--; + } + } while (aLftIndex <= aRghIndex); + + return aLftIndex; +} + +// ======================================================================= +// function : GetSubVolumes +// purpose : Arranges node triangles into bins +// ======================================================================= +void OpenGl_BinnedBVHBuilder::GetSubVolumes (OpenGl_RaytraceScene& theGeometry, + const OpenGl_BVHNode& theNode, + OpenGl_BinVector& theBins, + const int theAxis) +{ + const float aMin = theNode.MinPoint()[theAxis]; + const float aMax = theNode.MaxPoint()[theAxis]; + + const float aStep = (aMax - aMin) / THE_NUMBER_OF_BINS; + + for (int aTri = theNode.BegTriangle(); aTri <= theNode.EndTriangle(); ++aTri) + { + float aCenter = theGeometry.CenterAxis (aTri, theAxis); + int aBinIndex = (int )floorf ((aCenter - aMin) * ( 1.0f / aStep)); + if (aBinIndex < 0) + { + aBinIndex = 0; + } + else if (aBinIndex >= THE_NUMBER_OF_BINS) + { + aBinIndex = THE_NUMBER_OF_BINS - 1; + } + + theBins[aBinIndex].Count++; + theBins[aBinIndex].Volume.Combine (theGeometry.Box (aTri)); + } +} + +namespace OpenGl_Raytrace +{ + // ======================================================================= + // function : IsRaytracedElement + // purpose : Checks to see if the element contains ray-trace geometry + // ======================================================================= + Standard_Boolean IsRaytracedElement (const OpenGl_ElementNode* theNode) + { + if (TelParray == theNode->type) + { + OpenGl_PrimitiveArray* anArray = dynamic_cast< OpenGl_PrimitiveArray* > (theNode->elem); + return anArray->PArray()->type >= TelPolygonsArrayType; + } + return Standard_False; + } + + // ======================================================================= + // function : IsRaytracedGroup + // purpose : Checks to see if the group contains ray-trace geometry + // ======================================================================= + Standard_Boolean IsRaytracedGroup (const OpenGl_Group *theGroup) + { + const OpenGl_ElementNode* aNode; + for (aNode = theGroup->FirstNode(); aNode != NULL; aNode = aNode->next) + { + if (IsRaytracedElement (aNode)) + { + return Standard_True; + } + } + return Standard_False; + } + + // ======================================================================= + // function : IsRaytracedStructure + // purpose : Checks to see if the structure contains ray-trace geometry + // ======================================================================= + Standard_Boolean IsRaytracedStructure (const OpenGl_Structure *theStructure) + { + for (OpenGl_ListOfGroup::Iterator anItg (theStructure->Groups()); + anItg.More(); anItg.Next()) + { + if (anItg.Value()->IsRaytracable()) + return Standard_True; + } + for (OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures()); + anIts.More(); anIts.Next()) + { + if (IsRaytracedStructure (anIts.Value())) + return Standard_True; + } + return Standard_False; + } +} + +#endif diff --git a/src/OpenGl/OpenGl_SceneGeometry.hxx b/src/OpenGl/OpenGl_SceneGeometry.hxx new file mode 100644 index 0000000000..d01af3cbaf --- /dev/null +++ b/src/OpenGl/OpenGl_SceneGeometry.hxx @@ -0,0 +1,360 @@ +// Created on: 2013-08-27 +// Created by: Denis BOGOLEPOV +// Copyright (c) 2013 OPEN CASCADE SAS +// +// The content of this file is subject to the Open CASCADE Technology Public +// License Version 6.5 (the "License"). You may not use the content of this file +// except in compliance with the License. Please obtain a copy of the License +// at http://www.opencascade.org and read it completely before using this file. +// +// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its +// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France. +// +// The Original Code and all software distributed under the License is +// distributed on an "AS IS" basis, without warranty of any kind, and the +// Initial Developer hereby disclaims all such warranties, including without +// limitation, any warranties of merchantability, fitness for a particular +// purpose or non-infringement. Please see the License for the specific terms +// and conditions governing the rights and limitations under the License. + +#ifndef _OpenGl_SceneGeometry_Header +#define _OpenGl_SceneGeometry_Header + +#ifdef HAVE_OPENCL + +#include +#include +#include + +namespace OpenGl_Raytrace +{ + //! Checks to see if the group contains ray-trace geometry. + Standard_Boolean IsRaytracedGroup (const OpenGl_Group* theGroup); + + //! Checks to see if the element contains ray-trace geometry. + Standard_Boolean IsRaytracedElement (const OpenGl_ElementNode* theNode); + + //! Checks to see if the structure contains ray-trace geometry. + Standard_Boolean IsRaytracedStructure (const OpenGl_Structure* theStructure); +} + +//! Stores properties of surface material. +class OpenGl_RaytraceMaterial +{ +public: + + //! Ambient reflection coefficient. + OpenGl_RTVec4f Ambient; + + //! Diffuse reflection coefficient. + OpenGl_RTVec4f Diffuse; + + //! Glossy reflection coefficient. + OpenGl_RTVec4f Specular; + + //! Material emission. + OpenGl_RTVec4f Emission; + + //! Specular reflection coefficient. + OpenGl_RTVec4f Reflection; + + //! Specular refraction coefficient. + OpenGl_RTVec4f Refraction; + + //! Material transparency. + OpenGl_RTVec4f Transparency; + +public: + + //! Creates new default material. + OpenGl_RaytraceMaterial(); + + //! Creates new material with specified properties. + OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient, + const OpenGl_RTVec4f& theDiffuse, + const OpenGl_RTVec4f& theSpecular); + + //! Creates new material with specified properties. + OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient, + const OpenGl_RTVec4f& theDiffuse, + const OpenGl_RTVec4f& theSpecular, + const OpenGl_RTVec4f& theEmission, + const OpenGl_RTVec4f& theTranspar); + + //! Creates new material with specified properties. + OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient, + const OpenGl_RTVec4f& theDiffuse, + const OpenGl_RTVec4f& theSpecular, + const OpenGl_RTVec4f& theEmission, + const OpenGl_RTVec4f& theTranspar, + const OpenGl_RTVec4f& theReflection, + const OpenGl_RTVec4f& theRefraction); + + //! Returns packed (serialized) representation of material. + const float* Packed() { return reinterpret_cast (this); } +}; + +//! Stores properties of OpenGL light source. +class OpenGl_RaytraceLight +{ +public: + + //! 'Ambient' intensity. + OpenGl_RTVec4f Ambient; + + //! 'Diffuse' intensity. + OpenGl_RTVec4f Diffuse; + + //! Position of light source (in terms of OpenGL). + OpenGl_RTVec4f Position; + + +public: + + //! Creates new light source. + OpenGl_RaytraceLight (const OpenGl_RTVec4f& theAmbient); + + //! Creates new light source. + OpenGl_RaytraceLight (const OpenGl_RTVec4f& theDiffuse, + const OpenGl_RTVec4f& thePosition); + + //! Returns packed (serialized) representation of light source. + const float* Packed() { return reinterpret_cast (this); } +}; + +//! Stores scene geometry data. +struct OpenGl_RaytraceScene +{ + //! AABB of 3D scene. + OpenGl_AABB AABB; + + //! Array of vertex normals. + OpenGl_RTArray4f Normals; + + //! Array of vertex coordinates. + OpenGl_RTArray4f Vertices; + + //! Array of scene triangles. + OpenGl_RTArray4i Triangles; + + //! Array of 'front' material properties. + std::vector > Materials; + + //! Array of properties of light sources. + std::vector > LightSources; + + //! Clears all scene geometry and material data. + void Clear(); + + //! Returns AABB of specified triangle. + OpenGl_AABB Box (const int theTriangle) const; + + //! Returns centroid of specified triangle. + OpenGl_RTVec4f Center (const int theTriangle) const; + + //! Returns centroid coordinate for specified axis. + float CenterAxis (const int theTriangle, const int theAxis) const; +}; + +//! Stores parameters of BVH tree node. +class OpenGl_BVHNode +{ + friend class OpenGl_BVH; + +public: + + //! Creates new empty BVH node. + OpenGl_BVHNode(); + + //! Creates new BVH node with specified data. + OpenGl_BVHNode (const OpenGl_RTVec4f& theMinPoint, + const OpenGl_RTVec4f& theMaxPoint, + const OpenGl_RTVec4i& theDataRcrd); + + //! Creates new leaf BVH node with specified data. + OpenGl_BVHNode (const OpenGl_RTVec4f& theMinPoint, + const OpenGl_RTVec4f& theMaxPoint, + const int theBegTriangle, + const int theEndTriangle); + + //! Creates new leaf BVH node with specified data. + OpenGl_BVHNode (const OpenGl_AABB& theAABB, + const int theBegTriangle, + const int theEndTriangle); + + //! Returns minimum point of node's AABB. + OpenGl_RTVec4f& MinPoint() { return myMinPoint; } + //! Returns maximum point of node's AABB. + OpenGl_RTVec4f& MaxPoint() { return myMaxPoint; } + + //! Returns minimum point of node's AABB. + const OpenGl_RTVec4f& MinPoint() const { return myMinPoint; } + //! Returns maximum point of node's AABB. + const OpenGl_RTVec4f& MaxPoint() const { return myMaxPoint; } + + //! Returns index of left child of inner node. + int LeftChild() const { return myDataRcrd.y(); } + //! Sets index of left child of inner node. + void SetLeftChild (int theChild) { myDataRcrd.y() = theChild; } + + //! Returns index of right child of inner node. + int RightChild() const { return myDataRcrd.z(); } + //! Sets index of right child of inner node. + void SetRightChild (int theChild) { myDataRcrd.z() = theChild; } + + //! Returns index of begin triangle of leaf node. + int BegTriangle() const { return myDataRcrd.y(); } + //! Sets index of begin triangle of leaf node. + void SetBegTriangle (int theIndex) { myDataRcrd.y() = theIndex; } + + //! Returns index of end triangle of leaf node. + int EndTriangle() const { return myDataRcrd.z(); } + //! Sets index of end triangle of leaf node. + void SetEndTriangle (int theIndex) { myDataRcrd.z() = theIndex; } + + //! Returns level of the node in BVH tree. + int Level() const { return myDataRcrd.w(); } + //! Sets level of the node in BVH tree. + void SetLevel (int theLevel) { myDataRcrd.w() = theLevel; } + + //! Is node a leaf (outer)? + bool IsOuter() const { return myDataRcrd.x() == 1; } + + //! Sets node type to 'outer'. + void SetOuter() { myDataRcrd.x() = 1; } + //! Sets node type to 'inner'. + void SetInner() { myDataRcrd.x() = 0; } + +private: + + //! Minimum point of node's bounding box. + OpenGl_RTVec4f myMinPoint; + //! Maximum point of node's bounding box. + OpenGl_RTVec4f myMaxPoint; + + //! Data vector (stores data fields of the node). + OpenGl_RTVec4i myDataRcrd; +}; + +//! Stores parameters of BVH tree. +class OpenGl_BVH +{ +public: + + //! Removes all tree nodes. + void CleanUp(); + + //! Adds new node to the tree. + int PushNode (const OpenGl_BVHNode& theNode); + + //! Returns node with specified index. + OpenGl_BVHNode Node (const int theIndex) const; + + //! Replaces node with specified index by the new one. + void SetNode (const int theIndex, const OpenGl_BVHNode& theNode); + + //! Returns array of node min points. + OpenGl_RTArray4f& MinPointBuffer() { return myMinPointBuffer; } + //! Returns array of node max points. + OpenGl_RTArray4f& MaxPointBuffer() { return myMaxPointBuffer; } + //! Returns array of node data records. + OpenGl_RTArray4i& DataRcrdBuffer() { return myDataRcrdBuffer; } + +private: + + //! Array of min points of BVH nodes. + OpenGl_RTArray4f myMinPointBuffer; + //! Array of max points of BVH nodes. + OpenGl_RTArray4f myMaxPointBuffer; + //! Array of data vectors of BVH nodes. + OpenGl_RTArray4i myDataRcrdBuffer; +}; + +//! Stores parameters of single node bin (slice of AABB). +struct OpenGl_BVHBin +{ + //! Creates new node bin. + OpenGl_BVHBin(): Count (0) { } + + //! Number of primitives in the bin. + int Count; + + //! AABB of the bin. + OpenGl_AABB Volume; +}; + +//! Node building task. +struct OpenGl_BVHNodeTask +{ + //! Creates new node building task. + OpenGl_BVHNodeTask(); + + //! Creates new node building task. + OpenGl_BVHNodeTask (const int theNodeToBuild, + const int theBegTriangle, + const int theEndTriangle); + + //! Index of building tree node. + int NodeToBuild; + //! Index of start node triangle. + int BegTriangle; + //! Index of final node triangle. + int EndTriangle; +}; + +//! The array of bins of BVH tree node. +typedef std::vector > OpenGl_BinVector; + +//! Binned SAH-based BVH builder. +class OpenGl_BinnedBVHBuilder +{ +public: + + //! Creates new binned BVH builder. + OpenGl_BinnedBVHBuilder(); + + //! Releases binned BVH builder. + ~OpenGl_BinnedBVHBuilder(); + + //! Builds BVH tree using binned SAH algorithm. + void Build (OpenGl_RaytraceScene& theGeometry, const float theEpsilon = 1e-3f); + + //! Sets maximum tree depth. + void SetMaxDepth (const int theMaxDepth); + + //! Clears previously constructed BVH tree. + void CleanUp(); + + //! Return constructed BVH tree. + OpenGl_BVH& Tree() { return myTree; } + +private: + + //! Builds node using task info. + void BuildNode (OpenGl_RaytraceScene& theGeometry, const int theTask); + + //! Arranges node triangles into bins. + void GetSubVolumes (OpenGl_RaytraceScene& theGeometry, const OpenGl_BVHNode& theNode, + OpenGl_BinVector& theBins, const int theAxis); + + //! Splits node triangles into two intervals for child nodes. + int SplitTriangles (OpenGl_RaytraceScene& theGeometry, const int theFirst, const int theLast, + OpenGl_BVHNode& theNode, int theBin, const int theAxis); + +private: + + //! Queue of node building tasks. + std::vector myNodeTasksQueue; + + //! Builded BVH tree. + OpenGl_BVH myTree; + + //! Maximum depth of BVH tree. + int myMaxDepth; +}; + +#endif +#endif diff --git a/src/OpenGl/OpenGl_Structure.cxx b/src/OpenGl/OpenGl_Structure.cxx index 5f9c9d7e0b..42ade61c9e 100644 --- a/src/OpenGl/OpenGl_Structure.cxx +++ b/src/OpenGl/OpenGl_Structure.cxx @@ -17,6 +17,10 @@ // purpose or non-infringement. Please see the License for the specific terms // and conditions governing the rights and limitations under the License. +#ifdef HAVE_CONFIG_H + #include +#endif + #include #include @@ -144,6 +148,10 @@ OpenGl_Structure::OpenGl_Structure () myNamedStatus(0), myZLayer(0) { +#if HAVE_OPENCL + myIsRaytracable = Standard_False; + myModificationState = 0; +#endif } // ======================================================================= @@ -161,12 +169,21 @@ OpenGl_Structure::~OpenGl_Structure() // function : SetTransformation // purpose : // ======================================================================= -void OpenGl_Structure::SetTransformation(const float *AMatrix) +void OpenGl_Structure::SetTransformation (const float *theMatrix) { if (!myTransformation) + { myTransformation = new OpenGl_Matrix(); + } - matcpy( myTransformation->mat, AMatrix ); + matcpy (myTransformation->mat, theMatrix); + +#ifdef HAVE_OPENCL + if (myIsRaytracable) + { + UpdateStateWithAncestorStructures(); + } +#endif } // ======================================================================= @@ -208,6 +225,13 @@ void OpenGl_Structure::SetAspectFace (const CALL_DEF_CONTEXTFILLAREA& theAspect) myAspectFace = new OpenGl_AspectFace(); } myAspectFace->SetAspect (theAspect); + +#ifdef HAVE_OPENCL + if (myIsRaytracable) + { + UpdateStateWithAncestorStructures(); + } +#endif } // ======================================================================= @@ -249,7 +273,11 @@ void OpenGl_Structure::SetHighlightBox (const Handle(OpenGl_Context)& theGlCtx, } else { +#ifndef HAVE_OPENCL myHighlightBox = new OpenGl_Group(); +#else + myHighlightBox = new OpenGl_Group (this); +#endif } CALL_DEF_CONTEXTLINE aContextLine; @@ -306,29 +334,151 @@ void OpenGl_Structure::ClearHighlightColor (const Handle(OpenGl_Context)& theGlC myHighlightColor = NULL; } +// ======================================================================= +// function : SetNamedStatus +// purpose : +// ======================================================================= +void OpenGl_Structure::SetNamedStatus (const Standard_Integer aStatus) +{ + myNamedStatus = aStatus; + +#ifdef HAVE_OPENCL + if (myIsRaytracable) + { + UpdateStateWithAncestorStructures(); + } +#endif +} + +#ifdef HAVE_OPENCL + +// ======================================================================= +// function : RegisterAncestorStructure +// purpose : +// ======================================================================= +void OpenGl_Structure::RegisterAncestorStructure (const OpenGl_Structure* theStructure) const +{ + for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next()) + { + if (anIt.Value() == theStructure) + { + return; + } + } + + myAncestorStructures.Append (theStructure); +} + +// ======================================================================= +// function : UnregisterAncestorStructure +// purpose : +// ======================================================================= +void OpenGl_Structure::UnregisterAncestorStructure (const OpenGl_Structure* theStructure) const +{ + for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next()) + { + if (anIt.Value() == theStructure) + { + myAncestorStructures.Remove (anIt); + return; + } + } +} + +// ======================================================================= +// function : UpdateStateWithAncestorStructures +// purpose : +// ======================================================================= +void OpenGl_Structure::UpdateStateWithAncestorStructures() const +{ + myModificationState++; + + for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next()) + { + anIt.Value()->UpdateStateWithAncestorStructures(); + } +} + +// ======================================================================= +// function : UpdateRaytracableWithAncestorStructures +// purpose : +// ======================================================================= +void OpenGl_Structure::UpdateRaytracableWithAncestorStructures() const +{ + myIsRaytracable = OpenGl_Raytrace::IsRaytracedStructure (this); + + if (!myIsRaytracable) + { + for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next()) + { + anIt.Value()->UpdateRaytracableWithAncestorStructures(); + } + } +} + +// ======================================================================= +// function : SetRaytracableWithAncestorStructures +// purpose : +// ======================================================================= +void OpenGl_Structure::SetRaytracableWithAncestorStructures() const +{ + myIsRaytracable = Standard_True; + + for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next()) + { + if (!anIt.Value()->IsRaytracable()) + { + anIt.Value()->SetRaytracableWithAncestorStructures(); + } + } +} + +#endif + // ======================================================================= // function : Connect // purpose : // ======================================================================= -void OpenGl_Structure::Connect (const OpenGl_Structure *AStructure) +void OpenGl_Structure::Connect (const OpenGl_Structure *theStructure) { - Disconnect (AStructure); - myConnected.Append(AStructure); + Disconnect (theStructure); + myConnected.Append (theStructure); + +#ifdef HAVE_OPENCL + if (theStructure->IsRaytracable()) + { + UpdateStateWithAncestorStructures(); + SetRaytracableWithAncestorStructures(); + } + + theStructure->RegisterAncestorStructure (this); +#endif } // ======================================================================= // function : Disconnect // purpose : // ======================================================================= -void OpenGl_Structure::Disconnect (const OpenGl_Structure *AStructure) +void OpenGl_Structure::Disconnect (const OpenGl_Structure *theStructure) { - OpenGl_ListOfStructure::Iterator its(myConnected); + OpenGl_ListOfStructure::Iterator its (myConnected); while (its.More()) { // Check for the given structure - if (its.Value() == AStructure) + if (its.Value() == theStructure) { - myConnected.Remove(its); + myConnected.Remove (its); + +#ifdef HAVE_OPENCL + if (theStructure->IsRaytracable()) + { + UpdateStateWithAncestorStructures(); + UpdateRaytracableWithAncestorStructures(); + } + + theStructure->UnregisterAncestorStructure (this); +#endif + return; } its.Next(); @@ -339,10 +489,15 @@ void OpenGl_Structure::Disconnect (const OpenGl_Structure *AStructure) // function : AddGroup // purpose : // ======================================================================= -OpenGl_Group * OpenGl_Structure::AddGroup () +OpenGl_Group * OpenGl_Structure::AddGroup() { // Create new group - OpenGl_Group *g = new OpenGl_Group; +#ifndef HAVE_OPENCL + OpenGl_Group *g = new OpenGl_Group(); +#else + OpenGl_Group *g = new OpenGl_Group (this); +#endif + myGroups.Append(g); return g; } @@ -359,9 +514,18 @@ void OpenGl_Structure::RemoveGroup (const Handle(OpenGl_Context)& theGlCtx, // Check for the given group if (anIter.Value() == theGroup) { - // Delete object - OpenGl_Element::Destroy (theGlCtx, const_cast (anIter.ChangeValue())); myGroups.Remove (anIter); + +#ifdef HAVE_OPENCL + if (theGroup->IsRaytracable()) + { + UpdateStateWithAncestorStructures(); + UpdateRaytracableWithAncestorStructures(); + } +#endif + + // Delete object + OpenGl_Element::Destroy (theGlCtx, const_cast (theGroup)); return; } } @@ -373,13 +537,29 @@ void OpenGl_Structure::RemoveGroup (const Handle(OpenGl_Context)& theGlCtx, // ======================================================================= void OpenGl_Structure::Clear (const Handle(OpenGl_Context)& theGlCtx) { +#ifdef HAVE_OPENCL + Standard_Boolean aRaytracableGroupDeleted (Standard_False); +#endif + // Release groups for (OpenGl_ListOfGroup::Iterator anIter (myGroups); anIter.More(); anIter.Next()) { +#ifdef HAVE_OPENCL + aRaytracableGroupDeleted |= anIter.Value()->IsRaytracable(); +#endif + // Delete objects OpenGl_Element::Destroy (theGlCtx, const_cast (anIter.ChangeValue())); } myGroups.Clear(); + +#ifdef HAVE_OPENCL + if (aRaytracableGroupDeleted) + { + UpdateStateWithAncestorStructures(); + UpdateRaytracableWithAncestorStructures(); + } +#endif } // ======================================================================= diff --git a/src/OpenGl/OpenGl_Structure.hxx b/src/OpenGl/OpenGl_Structure.hxx index 8d49896ee1..994e6ad13e 100644 --- a/src/OpenGl/OpenGl_Structure.hxx +++ b/src/OpenGl/OpenGl_Structure.hxx @@ -30,6 +30,7 @@ #include #include +#include #include @@ -39,6 +40,7 @@ typedef NCollection_List OpenGl_ListOfStructure; class OpenGl_Structure : public OpenGl_Element { + friend class OpenGl_Group; public: @@ -65,7 +67,9 @@ public: void ClearHighlightColor (const Handle(OpenGl_Context)& theGlCtx); - void SetNamedStatus (const Standard_Integer aStatus) { myNamedStatus = aStatus; } + void SetNamedStatus (const Standard_Integer aStatus); + + Standard_Boolean IsVisible() const { return !(myNamedStatus & OPENGL_NS_HIDE); } void SetClipPlanes (const Graphic3d_SetOfHClipPlane& thePlanes) { myClipPlanes = thePlanes; } @@ -87,17 +91,64 @@ public: virtual void Release (const Handle(OpenGl_Context)& theGlCtx); //! This method releases GL resources without actual elements destruction. - //! As result structure could be correctly destroyed layter without GL context + //! As result structure could be correctly destroyed layer without GL context //! (after last window was closed for example). //! //! Notice however that reusage of this structure after calling this method is incorrect //! and will lead to broken visualization due to loosed data. void ReleaseGlResources (const Handle(OpenGl_Context)& theGlCtx); + //! Returns list of OpenGL groups. + const OpenGl_ListOfGroup& Groups() const { return myGroups; } + + //! Returns list of connected OpenGL structures. + const OpenGl_ListOfStructure& ConnectedStructures() const { return myConnected; } + + //! Returns OpenGL face aspect. + const OpenGl_AspectFace* AspectFace() const { return myAspectFace; } + + //! Returns OpenGL transformation matrix. + const OpenGl_Matrix* Transformation() const { return myTransformation; } + + //! Returns OpenGL persistent translation. + const TEL_TRANSFORM_PERSISTENCE* PersistentTranslation() const { return myTransPers; } + +#ifdef HAVE_OPENCL + + //! Returns structure modification state (for ray-tracing). + Standard_Size ModificationState() const { return myModificationState; } + + //! Resets structure modification state (for ray-tracing) + void ResetModificationState() const { myModificationState = 0; } + + //! Is the structure ray-tracable (contains ray-tracable elements)? + Standard_Boolean IsRaytracable() const { return myIsRaytracable; } + +#endif + protected: virtual ~OpenGl_Structure(); +#ifdef HAVE_OPENCL + + //! Registers ancestor connected structure (for updating ray-tracing state). + void RegisterAncestorStructure (const OpenGl_Structure* theStructure) const; + + //! Unregisters ancestor connected structure (for updating ray-tracing state). + void UnregisterAncestorStructure (const OpenGl_Structure* theStructure) const; + + //! Updates modification state for structure and its parents. + void UpdateStateWithAncestorStructures() const; + + //! Updates ray-tracable status for structure and its parents. + void UpdateRaytracableWithAncestorStructures() const; + + //! Sets ray-tracable status for structure and its parents. + void SetRaytracableWithAncestorStructures() const; + +#endif + protected: //Structure_LABBegin @@ -119,6 +170,12 @@ protected: OpenGl_ListOfGroup myGroups; Graphic3d_SetOfHClipPlane myClipPlanes; +#ifdef HAVE_OPENCL + mutable OpenGl_ListOfStructure myAncestorStructures; + mutable Standard_Boolean myIsRaytracable; + mutable Standard_Size myModificationState; +#endif + public: DEFINE_STANDARD_ALLOC diff --git a/src/OpenGl/OpenGl_View.cxx b/src/OpenGl/OpenGl_View.cxx index 32e4b2d58f..29244efba8 100644 --- a/src/OpenGl/OpenGl_View.cxx +++ b/src/OpenGl/OpenGl_View.cxx @@ -17,6 +17,11 @@ // purpose or non-infringement. Please see the License for the specific terms // and conditions governing the rights and limitations under the License. + +#ifdef HAVE_CONFIG_H + #include +#endif + #include #include @@ -123,6 +128,10 @@ OpenGl_View::OpenGl_View (const CALL_DEF_VIEWCONTEXT &AContext) myIntShadingMethod = TEL_SM_FLAT; break; } + +#ifdef HAVE_OPENCL + myModificationState = 1; // initial state +#endif } /*----------------------------------------------------------------------*/ @@ -165,8 +174,21 @@ void OpenGl_View::SetTextureEnv (const Handle(OpenGl_Context)& theCtx, myTextureEnv = new OpenGl_Texture (theTexture->GetParams()); Handle(Image_PixMap) anImage = theTexture->GetImage(); - if( !anImage.IsNull()) + if (!anImage.IsNull()) myTextureEnv->Init (theCtx, *anImage.operator->(), theTexture->Type()); + +#ifdef HAVE_OPENCL + myModificationState++; +#endif +} + +void OpenGl_View::SetSurfaceDetail (const Visual3d_TypeOfSurfaceDetail theMode) +{ + mySurfaceDetail = theMode; + +#ifdef HAVE_OPENCL + myModificationState++; +#endif } /*----------------------------------------------------------------------*/ diff --git a/src/OpenGl/OpenGl_View.hxx b/src/OpenGl/OpenGl_View.hxx index fb1636dbb4..6419c1cba4 100644 --- a/src/OpenGl/OpenGl_View.hxx +++ b/src/OpenGl/OpenGl_View.hxx @@ -109,7 +109,7 @@ class OpenGl_View : public MMgt_TShared void SetTextureEnv (const Handle(OpenGl_Context)& theCtx, const Handle(Graphic3d_TextureEnv)& theTexture); - void SetSurfaceDetail (const Visual3d_TypeOfSurfaceDetail AMode) { mySurfaceDetail = AMode; } + void SetSurfaceDetail (const Visual3d_TypeOfSurfaceDetail AMode); void SetBackfacing (const Standard_Integer AMode); void SetLights (const CALL_DEF_VIEWCONTEXT &AContext); void SetAntiAliasing (const Standard_Boolean AMode) { myAntiAliasing = AMode; } @@ -189,6 +189,26 @@ class OpenGl_View : public MMgt_TShared const Aspect_CLayer2d& theCUnderLayer, const Aspect_CLayer2d& theCOverLayer); + + void DrawBackground (const Handle(OpenGl_Workspace)& theWorkspace); + + //! Returns list of OpenGL Z-layers. + const OpenGl_LayerList& LayerList() const { return myZLayers; } + + //! Returns list of openGL light sources. + const OpenGl_ListOfLight& LightList() const { return myLights; } + + //! Returns OpenGL environment map. + const Handle(OpenGl_Texture)& TextureEnv() const { return myTextureEnv; } + + //! Returns visualization mode for objects in the view. + Visual3d_TypeOfSurfaceDetail SurfaceDetail() const { return mySurfaceDetail; } + +#ifdef HAVE_OPENCL + //! Returns modification state for ray-tracing. + Standard_Size ModificationState() const { return myModificationState; } +#endif + public: DEFINE_STANDARD_RTTI(OpenGl_View) // Type definition @@ -250,6 +270,10 @@ public: Standard_Boolean myViewMappingChanged; Standard_Boolean myLightSourcesChanged; +#ifdef HAVE_OPENCL + Standard_Size myModificationState; +#endif + public: DEFINE_STANDARD_ALLOC }; diff --git a/src/OpenGl/OpenGl_View_2.cxx b/src/OpenGl/OpenGl_View_2.cxx index 7200c3a040..c11b1f9cc3 100644 --- a/src/OpenGl/OpenGl_View_2.cxx +++ b/src/OpenGl/OpenGl_View_2.cxx @@ -636,73 +636,14 @@ call_util_mat_mul( matrix3 mat_a, matrix3 mat_b, matrix3 mat_c) /*----------------------------------------------------------------------*/ -//call_func_redraw_all_structs_proc -void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext, - const Handle(OpenGl_Workspace) &AWorkspace, - const Graphic3d_CView& ACView, - const Aspect_CLayer2d& ACUnderLayer, - const Aspect_CLayer2d& ACOverLayer) +void OpenGl_View::DrawBackground (const Handle(OpenGl_Workspace) &AWorkspace) { - // Store and disable current clipping planes - const Handle(OpenGl_Context)& aContext = AWorkspace->GetGlContext(); - const Standard_Integer aMaxClipPlanes = aContext->MaxClipPlanes(); - const GLenum lastid = GL_CLIP_PLANE0 + aMaxClipPlanes; - OPENGL_CLIP_PLANE *oldPlanes = new OPENGL_CLIP_PLANE[aMaxClipPlanes]; - OPENGL_CLIP_PLANE *ptrPlane = oldPlanes; - GLenum planeid = GL_CLIP_PLANE0; - for ( ; planeid < lastid; planeid++, ptrPlane++ ) - { - glGetClipPlane( planeid, ptrPlane->Equation ); - if ( ptrPlane->isEnabled ) - { - glDisable( planeid ); - ptrPlane->isEnabled = GL_TRUE; - } - else - { - ptrPlane->isEnabled = GL_FALSE; - } - } - - // Set OCCT state uniform variables - - if (!aContext->ShaderManager()->IsEmpty()) - { - if (myLightSourcesChanged) - { - aContext->ShaderManager()->UpdateLightSourceStateTo (&myLights); - myLightSourcesChanged = Standard_False; - } - - if (myViewMappingChanged) - { - aContext->ShaderManager()->UpdateProjectionStateTo (myMappingMatrix); - myViewMappingChanged = Standard_False; - } - - if (myOrientationChanged) - { - aContext->ShaderManager()->UpdateWorldViewStateTo (myOrientationMatrix); - myOrientationChanged = Standard_False; - } - - if (aContext->ShaderManager()->ModelWorldState().Index() == 0) - { - Tmatrix3 aModelWorldState = { { 1.f, 0.f, 0.f, 0.f }, - { 0.f, 1.f, 0.f, 0.f }, - { 0.f, 0.f, 1.f, 0.f }, - { 0.f, 0.f, 0.f, 1.f } }; - - aContext->ShaderManager()->UpdateModelWorldStateTo (aModelWorldState); - } - } - ///////////////////////////////////////////////////////////////////////////// // Step 1: Prepare for redraw // Render background if ( (AWorkspace->NamedStatus & OPENGL_NS_WHITEBACK) == 0 && - ( myBgTexture.TexId != 0 || myBgGradient.type != Aspect_GFM_NONE ) ) + ( myBgTexture.TexId != 0 || myBgGradient.type != Aspect_GFM_NONE ) ) { const Standard_Integer aViewWidth = AWorkspace->Width(); const Standard_Integer aViewHeight = AWorkspace->Height(); @@ -723,8 +664,8 @@ void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext, // - gradient fill type is not Aspect_GFM_NONE and // - either background texture is no specified or it is drawn in Aspect_FM_CENTERED mode if ( ( myBgGradient.type != Aspect_GFM_NONE ) && - ( myBgTexture.TexId == 0 || myBgTexture.Style == Aspect_FM_CENTERED || - myBgTexture.Style == Aspect_FM_NONE ) ) + ( myBgTexture.TexId == 0 || myBgTexture.Style == Aspect_FM_CENTERED || + myBgTexture.Style == Aspect_FM_NONE ) ) { Tfloat* corner1 = 0;/* -1,-1*/ Tfloat* corner2 = 0;/* 1,-1*/ @@ -735,63 +676,63 @@ void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext, switch( myBgGradient.type ) { - case Aspect_GFM_HOR: - corner1 = myBgGradient.color1.rgb; - corner2 = myBgGradient.color2.rgb; - corner3 = myBgGradient.color2.rgb; - corner4 = myBgGradient.color1.rgb; - break; - case Aspect_GFM_VER: - corner1 = myBgGradient.color2.rgb; - corner2 = myBgGradient.color2.rgb; - corner3 = myBgGradient.color1.rgb; - corner4 = myBgGradient.color1.rgb; - break; - case Aspect_GFM_DIAG1: - corner2 = myBgGradient.color2.rgb; - corner4 = myBgGradient.color1.rgb; - dcorner1 [0] = dcorner2[0] = 0.5F * (corner2[0] + corner4[0]); - dcorner1 [1] = dcorner2[1] = 0.5F * (corner2[1] + corner4[1]); - dcorner1 [2] = dcorner2[2] = 0.5F * (corner2[2] + corner4[2]); - corner1 = dcorner1; - corner3 = dcorner2; - break; - case Aspect_GFM_DIAG2: - corner1 = myBgGradient.color2.rgb; - corner3 = myBgGradient.color1.rgb; - dcorner1 [0] = dcorner2[0] = 0.5F * (corner1[0] + corner3[0]); - dcorner1 [1] = dcorner2[1] = 0.5F * (corner1[1] + corner3[1]); - dcorner1 [2] = dcorner2[2] = 0.5F * (corner1[2] + corner3[2]); - corner2 = dcorner1; - corner4 = dcorner2; - break; - case Aspect_GFM_CORNER1: - corner1 = myBgGradient.color2.rgb; - corner2 = myBgGradient.color2.rgb; - corner3 = myBgGradient.color2.rgb; - corner4 = myBgGradient.color1.rgb; - break; - case Aspect_GFM_CORNER2: - corner1 = myBgGradient.color2.rgb; - corner2 = myBgGradient.color2.rgb; - corner3 = myBgGradient.color1.rgb; - corner4 = myBgGradient.color2.rgb; - break; - case Aspect_GFM_CORNER3: - corner1 = myBgGradient.color2.rgb; - corner2 = myBgGradient.color1.rgb; - corner3 = myBgGradient.color2.rgb; - corner4 = myBgGradient.color2.rgb; - break; - case Aspect_GFM_CORNER4: - corner1 = myBgGradient.color1.rgb; - corner2 = myBgGradient.color2.rgb; - corner3 = myBgGradient.color2.rgb; - corner4 = myBgGradient.color2.rgb; - break; - default: - //printf("gradient background type not right\n"); - break; + case Aspect_GFM_HOR: + corner1 = myBgGradient.color1.rgb; + corner2 = myBgGradient.color2.rgb; + corner3 = myBgGradient.color2.rgb; + corner4 = myBgGradient.color1.rgb; + break; + case Aspect_GFM_VER: + corner1 = myBgGradient.color2.rgb; + corner2 = myBgGradient.color2.rgb; + corner3 = myBgGradient.color1.rgb; + corner4 = myBgGradient.color1.rgb; + break; + case Aspect_GFM_DIAG1: + corner2 = myBgGradient.color2.rgb; + corner4 = myBgGradient.color1.rgb; + dcorner1 [0] = dcorner2[0] = 0.5F * (corner2[0] + corner4[0]); + dcorner1 [1] = dcorner2[1] = 0.5F * (corner2[1] + corner4[1]); + dcorner1 [2] = dcorner2[2] = 0.5F * (corner2[2] + corner4[2]); + corner1 = dcorner1; + corner3 = dcorner2; + break; + case Aspect_GFM_DIAG2: + corner1 = myBgGradient.color2.rgb; + corner3 = myBgGradient.color1.rgb; + dcorner1 [0] = dcorner2[0] = 0.5F * (corner1[0] + corner3[0]); + dcorner1 [1] = dcorner2[1] = 0.5F * (corner1[1] + corner3[1]); + dcorner1 [2] = dcorner2[2] = 0.5F * (corner1[2] + corner3[2]); + corner2 = dcorner1; + corner4 = dcorner2; + break; + case Aspect_GFM_CORNER1: + corner1 = myBgGradient.color2.rgb; + corner2 = myBgGradient.color2.rgb; + corner3 = myBgGradient.color2.rgb; + corner4 = myBgGradient.color1.rgb; + break; + case Aspect_GFM_CORNER2: + corner1 = myBgGradient.color2.rgb; + corner2 = myBgGradient.color2.rgb; + corner3 = myBgGradient.color1.rgb; + corner4 = myBgGradient.color2.rgb; + break; + case Aspect_GFM_CORNER3: + corner1 = myBgGradient.color2.rgb; + corner2 = myBgGradient.color1.rgb; + corner3 = myBgGradient.color2.rgb; + corner4 = myBgGradient.color2.rgb; + break; + case Aspect_GFM_CORNER4: + corner1 = myBgGradient.color1.rgb; + corner2 = myBgGradient.color2.rgb; + corner3 = myBgGradient.color2.rgb; + corner4 = myBgGradient.color2.rgb; + break; + default: + //printf("gradient background type not right\n"); + break; } // Save GL parameters @@ -831,7 +772,7 @@ void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext, GLfloat texX_range = 1.F; // texture coordinate GLfloat texY_range = 1.F; // texture coordinate - // Set up for stretching or tiling + // Set up for stretching or tiling GLfloat x_offset, y_offset; if ( myBgTexture.Style == Aspect_FM_CENTERED ) { @@ -891,6 +832,76 @@ void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext, else glDisable (GL_DITHER); } +} + +/*----------------------------------------------------------------------*/ + +//call_func_redraw_all_structs_proc +void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext, + const Handle(OpenGl_Workspace) &AWorkspace, + const Graphic3d_CView& ACView, + const Aspect_CLayer2d& ACUnderLayer, + const Aspect_CLayer2d& ACOverLayer) +{ + // Store and disable current clipping planes + const Handle(OpenGl_Context)& aContext = AWorkspace->GetGlContext(); + const Standard_Integer aMaxClipPlanes = aContext->MaxClipPlanes(); + const GLenum lastid = GL_CLIP_PLANE0 + aMaxClipPlanes; + OPENGL_CLIP_PLANE *oldPlanes = new OPENGL_CLIP_PLANE[aMaxClipPlanes]; + OPENGL_CLIP_PLANE *ptrPlane = oldPlanes; + GLenum planeid = GL_CLIP_PLANE0; + for ( ; planeid < lastid; planeid++, ptrPlane++ ) + { + glGetClipPlane( planeid, ptrPlane->Equation ); + if ( ptrPlane->isEnabled ) + { + glDisable( planeid ); + ptrPlane->isEnabled = GL_TRUE; + } + else + { + ptrPlane->isEnabled = GL_FALSE; + } + } + + // Set OCCT state uniform variables + + if (!aContext->ShaderManager()->IsEmpty()) + { + if (myLightSourcesChanged) + { + aContext->ShaderManager()->UpdateLightSourceStateTo (&myLights); + myLightSourcesChanged = Standard_False; + } + + if (myViewMappingChanged) + { + aContext->ShaderManager()->UpdateProjectionStateTo (myMappingMatrix); + myViewMappingChanged = Standard_False; + } + + if (myOrientationChanged) + { + aContext->ShaderManager()->UpdateWorldViewStateTo (myOrientationMatrix); + myOrientationChanged = Standard_False; + } + + if (aContext->ShaderManager()->ModelWorldState().Index() == 0) + { + Tmatrix3 aModelWorldState = { { 1.f, 0.f, 0.f, 0.f }, + { 0.f, 1.f, 0.f, 0.f }, + { 0.f, 0.f, 1.f, 0.f }, + { 0.f, 0.f, 0.f, 1.f } }; + + aContext->ShaderManager()->UpdateModelWorldStateTo (aModelWorldState); + } + } + + ///////////////////////////////////////////////////////////////////////////// + // Step 1: Prepare for redraw + + // Render background + DrawBackground (AWorkspace); // Switch off lighting by default glDisable(GL_LIGHTING); @@ -1028,8 +1039,6 @@ D = -[Px,Py,Pz] dot |Nx| // Apply clipping planes { - const Handle(OpenGl_Context)& aContext = AWorkspace->GetGlContext(); - if (myZClip.Back.IsOn || myZClip.Front.IsOn) { const GLdouble ramp = myExtra.map.fpd - myExtra.map.bpd; diff --git a/src/OpenGl/OpenGl_Workspace.cxx b/src/OpenGl/OpenGl_Workspace.cxx index e7562f94b9..e6fdf678d6 100644 --- a/src/OpenGl/OpenGl_Workspace.cxx +++ b/src/OpenGl/OpenGl_Workspace.cxx @@ -17,6 +17,11 @@ // purpose or non-infringement. Please see the License for the specific terms // and conditions governing the rights and limitations under the License. + +#ifdef HAVE_CONFIG_H + #include +#endif + #include #include @@ -28,6 +33,7 @@ #include #include #include +#include #include #include @@ -116,6 +122,21 @@ OpenGl_Workspace::OpenGl_Workspace (const Handle(OpenGl_Display)& theDisplay, // Polygon Offset EnablePolygonOffset(); + +#ifdef HAVE_OPENCL + + myComputeInitStatus = OpenGl_CLIS_NONE; + + myViewModificationStatus = 0; + myLayersModificationStatus = 0; + + myRaytraceOutputTexture[0] = 0; + myRaytraceOutputTexture[1] = 0; + + myIsRaytraceDataValid = Standard_False; + myToUpdateRaytraceData = Standard_False; + +#endif } // ======================================================================= @@ -135,6 +156,9 @@ Standard_Boolean OpenGl_Workspace::SetImmediateModeDrawToFront (const Standard_B // ======================================================================= OpenGl_Workspace::~OpenGl_Workspace() { +#ifdef HAVE_OPENCL + ReleaseOpenCL(); +#endif } // ======================================================================= @@ -483,11 +507,29 @@ void OpenGl_Workspace::Redraw (const Graphic3d_CView& theCView, toSwap = 0; // no need to swap buffers } - Redraw1 (theCView, theCUnderLayer, theCOverLayer, toSwap); - if (aFrameBuffer == NULL || !myTransientDrawToFront) +#ifdef HAVE_OPENCL + if (!theCView.IsRaytracing || myComputeInitStatus == OpenGl_CLIS_FAIL) + { +#endif + Redraw1 (theCView, theCUnderLayer, theCOverLayer, toSwap); + if (aFrameBuffer == NULL || !myTransientDrawToFront) + { + RedrawImmediatMode(); + } + + theCView.WasRedrawnGL = Standard_True; +#ifdef HAVE_OPENCL + } + else { - RedrawImmediatMode(); + int aSizeX = aFrameBuffer != NULL ? aFrameBuffer->GetVPSizeX() : myWidth; + int aSizeY = aFrameBuffer != NULL ? aFrameBuffer->GetVPSizeY() : myHeight; + + Raytrace (theCView, aSizeX, aSizeY, toSwap); + + theCView.WasRedrawnGL = Standard_False; } +#endif if (aFrameBuffer != NULL) { diff --git a/src/OpenGl/OpenGl_Workspace.hxx b/src/OpenGl/OpenGl_Workspace.hxx index c17a0064f6..4f13dbda6b 100644 --- a/src/OpenGl/OpenGl_Workspace.hxx +++ b/src/OpenGl/OpenGl_Workspace.hxx @@ -20,6 +20,13 @@ #ifndef _OpenGl_Workspace_Header #define _OpenGl_Workspace_Header +#ifdef HAVE_OPENCL + #include + #include + + #include +#endif + #include #include @@ -48,6 +55,9 @@ #include #include #include +#ifdef HAVE_OPENCL + #include +#endif #include #include @@ -62,8 +72,8 @@ class OpenGl_Structure; class OpenGl_Element; class Image_PixMap; -//! Reprepsents window with GL context. -//! Provides methods to render primitives and maintan GL state. +//! Represents window with GL context. +//! Provides methods to render primitives and maintain GL state. class OpenGl_Workspace : public OpenGl_Window { public: @@ -208,6 +218,171 @@ protected: void setTextureParams (Handle(OpenGl_Texture)& theTexture, const Handle(Graphic3d_TextureParams)& theParams); +#ifdef HAVE_OPENCL + +public: + + //! Returns information about OpenCL device used for computations. + Standard_Boolean GetOpenClDeviceInfo ( + NCollection_DataMap& theInfo) const; + +protected: + + //! Describes result of OpenCL initializing. + enum OpenClInitStatus + { + OpenGl_CLIS_NONE, + OpenGl_CLIS_INIT, + OpenGl_CLIS_FAIL + }; + +protected: //! @name methods related to ray-tracing + + //! Updates 3D scene geometry for ray-tracing. + Standard_Boolean UpdateRaytraceGeometry (Standard_Boolean theCheck); + + //! Checks to see if the structure is modified. + Standard_Boolean CheckRaytraceStructure (const OpenGl_Structure* theStructure); + + //! Updates 3D scene light sources for ray-tracing. + Standard_Boolean UpdateRaytraceLightSources (const GLdouble theInvModelView[16]); + + //! Updates environment map for ray-tracing. + Standard_Boolean UpdateRaytraceEnvironmentMap(); + + //! Adds OpenGL structure to ray-traced scene geometry. + Standard_Boolean AddRaytraceStructure (const OpenGl_Structure* theStruct, + const float* theTrans, std::set& theElements); + + //! Adds OpenGL primitive array to ray-traced scene geometry. + Standard_Boolean AddRaytracePrimitiveArray ( + const CALL_DEF_PARRAY* theArray, int theMatID, const float* theTrans); + + //! Adds vertex indices from OpenGL primitive array to ray-traced scene geometry. + Standard_Boolean AddRaytraceVertexIndices (const CALL_DEF_PARRAY* theArray, + int theFirstVert, int theVertOffset, int theVertNum, int theMatID); + + //! Adds OpenGL triangle array to ray-traced scene geometry. + Standard_Boolean AddRaytraceTriangleArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, int theVertOffset, int theVertNum, int theMatID); + + //! Adds OpenGL triangle fan array to ray-traced scene geometry. + Standard_Boolean AddRaytraceTriangleFanArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, int theVertOffset, int theVertNum, int theMatID); + + //! Adds OpenGL triangle fan array to ray-traced scene geometry. + Standard_Boolean AddRaytraceTriangleStripArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, int theVertOffset, int theVertNum, int theMatID); + + //! Adds OpenGL quadrangle array to ray-traced scene geometry. + Standard_Boolean AddRaytraceQuadrangleArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, int theVertOffset, int theVertNum, int theMatID); + + //! Adds OpenGL quadrangle strip array to ray-traced scene geometry. + Standard_Boolean AddRaytraceQuadrangleStripArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, int theVertOffset, int theVertNum, int theMatID); + + //! Adds OpenGL polygon array to ray-traced scene geometry. + Standard_Boolean AddRaytracePolygonArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, int theVertOffset, int theVertNum, int theMatID); + + //! Initializes OpenCL resources. + Standard_Boolean InitOpenCL(); + + //! Releases OpenCL resources. + void ReleaseOpenCL(); + + //! Resizes OpenCL output image. + Standard_Boolean ResizeRaytraceOutputBuffer (const cl_int theSizeX, const cl_int theSizeY); + + //! Writes scene geometry to OpenCl device. + Standard_Boolean WriteRaytraceSceneToDevice(); + + //! Runs OpenCL ray-tracing kernels. + Standard_Boolean RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView, + const GLfloat theOrigins[16], + const GLfloat theDirects[16], + const int theSizeX, + const int theSizeY); + + //! Redraws the window using OpenCL ray tracing. + Standard_Boolean Raytrace (const Graphic3d_CView& theCView, + const int theSizeX, int theSizeY, const Tint theToSwap); + +protected: //! @name fields related to ray-tracing + + //! Result of OpenCL initialization. + OpenClInitStatus myComputeInitStatus; + //! Is ATI/AMD OpenCL platform used? + Standard_Boolean myIsAmdComputePlatform; + + //! Is geometry data valid? + Standard_Boolean myIsRaytraceDataValid; + //! Is geometry data musty be updated? + Standard_Boolean myToUpdateRaytraceData; + //! 3D scene geometry data for ray-tracing. + OpenGl_RaytraceScene myRaytraceSceneData; + + //! Radius of bounding sphere of the scene. + float myRaytraceSceneRadius; + //! Scene epsilon to prevent self-intersections. + float myRaytraceSceneEpsilon; + + //! Binned SAH-based BVH builder. + OpenGl_BinnedBVHBuilder myBVHBuilder; + + //! OpenCL context. + cl_context myComputeContext; + //! OpenCL command queue. + cl_command_queue myRaytraceQueue; + //! OpenCL computing program. + cl_program myRaytraceProgram; + //! OpenCL ray-tracing render kernel. + cl_kernel myRaytraceRenderKernel; + //! OpenCL adaptive anti-aliasing kernel. + cl_kernel myRaytraceSmoothKernel; + + //! OpenCL image to store environment map. + cl_mem myRaytraceEnvironment; + //! OpenCL image to store rendering result. + cl_mem myRaytraceOutputImage; + //! OpenCL image to store anti-aliasing result. + cl_mem myRaytraceOutputImageSmooth; + + //! OpenGL output texture handle. + GLuint myRaytraceOutputTexture[2]; + + //! OpenCL buffer of vertex normals. + cl_mem myRaytraceNormalBuffer; + //! OpenCL buffer of vertex coordinates. + cl_mem myRaytraceVertexBuffer; + //! OpenCL buffer of indices of triangle vertices. + cl_mem myRaytraceTriangleBuffer; + + //! OpenCL buffer of minimum points of BVH nodes. + cl_mem myRaytraceNodeMinPointBuffer; + //! OpenCL buffer of maximum points of BVH nodes. + cl_mem myRaytraceNodeMaxPointBuffer; + //! OpenCL buffer of data records of BVH nodes. + cl_mem myRaytraceNodeDataRcrdBuffer; + + //! OpenCL buffer of material properties. + cl_mem myRaytraceMaterialBuffer; + + //! OpenCL buffer of light source properties. + cl_mem myRaytraceLightSourceBuffer; + + //! State of OpenGL view. + Standard_Size myViewModificationStatus; + + //! State of OpenGL layer list. + Standard_Size myLayersModificationStatus; + + //! State of OpenGL elements reflected to ray-tracing. + std::map myStructureStates; + +#endif // HAVE_OPENCL + protected: //! @name protected fields Handle(OpenGl_PrinterContext) myPrintContext; diff --git a/src/OpenGl/OpenGl_Workspace_Raytrace.cxx b/src/OpenGl/OpenGl_Workspace_Raytrace.cxx new file mode 100644 index 0000000000..534c7c3ec0 --- /dev/null +++ b/src/OpenGl/OpenGl_Workspace_Raytrace.cxx @@ -0,0 +1,2176 @@ +// Created on: 2013-08-27 +// Created by: Denis BOGOLEPOV +// Copyright (c) 2013 OPEN CASCADE SAS +// +// The content of this file is subject to the Open CASCADE Technology Public +// License Version 6.5 (the "License"). You may not use the content of this file +// except in compliance with the License. Please obtain a copy of the License +// at http://www.opencascade.org and read it completely before using this file. +// +// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its +// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France. +// +// The Original Code and all software distributed under the License is +// distributed on an "AS IS" basis, without warranty of any kind, and the +// Initial Developer hereby disclaims all such warranties, including without +// limitation, any warranties of merchantability, fitness for a particular +// purpose or non-infringement. Please see the License for the specific terms +// and conditions governing the rights and limitations under the License. + +#ifdef HAVE_CONFIG_H + #include +#endif + +#ifdef HAVE_OPENCL + +#if defined(_WIN32) + + #include + #include + + #pragma comment (lib, "DelayImp.lib") + #pragma comment (lib, "OpenCL.lib") + +#elif defined(__APPLE__) && !defined(MACOSX_USE_GLX) + #include +#else + #include +#endif + +#include +#include +#include +#include + +using namespace OpenGl_Raytrace; + +//! Use this macro to output ray-tracing debug info +// #define RAY_TRACE_PRINT_INFO + +#ifdef RAY_TRACE_PRINT_INFO + #include +#endif + +//! OpenCL source of ray-tracing kernels. +extern const char THE_RAY_TRACE_OPENCL_SOURCE[]; + +// ======================================================================= +// function : MatVecMult +// purpose : Multiples 4x4 matrix by 4D vector +// ======================================================================= +template< typename T > +OpenGl_RTVec4f MatVecMult (const T m[16], const OpenGl_RTVec4f& v) +{ + return OpenGl_RTVec4f ( + static_cast (m[ 0] * v.x() + m[ 4] * v.y() + + m[ 8] * v.z() + m[12] * v.w()), + static_cast (m[ 1] * v.x() + m[ 5] * v.y() + + m[ 9] * v.z() + m[13] * v.w()), + static_cast (m[ 2] * v.x() + m[ 6] * v.y() + + m[10] * v.z() + m[14] * v.w()), + static_cast (m[ 3] * v.x() + m[ 7] * v.y() + + m[11] * v.z() + m[15] * v.w())); +} + +// ======================================================================= +// function : UpdateRaytraceEnvironmentMap +// purpose : Updates environment map for ray-tracing +// ======================================================================= +Standard_Boolean OpenGl_Workspace::UpdateRaytraceEnvironmentMap() +{ + if (myView.IsNull()) + return Standard_False; + + if (myViewModificationStatus == myView->ModificationState()) + return Standard_True; + + cl_int anError = CL_SUCCESS; + + if (myRaytraceEnvironment != NULL) + clReleaseMemObject (myRaytraceEnvironment); + + int aSizeX = 1; + int aSizeY = 1; + + if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE) + { + aSizeX = (myView->TextureEnv()->SizeX() <= 0) ? 1 : myView->TextureEnv()->SizeX(); + aSizeY = (myView->TextureEnv()->SizeY() <= 0) ? 1 : myView->TextureEnv()->SizeY(); + } + + cl_image_format aImageFormat; + + aImageFormat.image_channel_order = CL_RGBA; + aImageFormat.image_channel_data_type = CL_FLOAT; + + myRaytraceEnvironment = clCreateImage2D (myComputeContext, CL_MEM_READ_ONLY, + &aImageFormat, aSizeX, aSizeY, 0, + NULL, &anError); + + cl_float* aPixelData = new cl_float[aSizeX * aSizeY * 4]; + + // Note: texture format is not compatible with OpenCL image + // (it's not possible to create image directly from texture) + + if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE) + { + myView->TextureEnv()->Bind (GetGlContext()); + + glGetTexImage (GL_TEXTURE_2D, + 0, + GL_RGBA, + GL_FLOAT, + aPixelData); + + myView->TextureEnv()->Unbind (GetGlContext()); + } + else + { + for (int aPixel = 0; aPixel < aSizeX * aSizeY * 4; ++aPixel) + aPixelData[aPixel] = 0.f; + } + + size_t anImageOffset[] = { 0, + 0, + 0 }; + + size_t anImageRegion[] = { aSizeX, + aSizeY, + 1 }; + + anError |= clEnqueueWriteImage (myRaytraceQueue, myRaytraceEnvironment, CL_TRUE, + anImageOffset, anImageRegion, 0, 0, aPixelData, + 0, NULL, NULL); +#ifdef RAY_TRACE_PRINT_INFO + if (anError != CL_SUCCESS) + std::cout << "Error! Failed to write environment map image!" << std::endl; +#endif + + delete[] aPixelData; + + myViewModificationStatus = myView->ModificationState(); + + return (anError == CL_SUCCESS); +} + +// ======================================================================= +// function : UpdateRaytraceGeometry +// purpose : Updates 3D scene geometry for ray tracing +// ======================================================================= +Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theCheck) +{ + if (myView.IsNull()) + return Standard_False; + + // Note: In 'check' mode the scene geometry is analyzed for modifications + // This is light-weight procedure performed for each frame + + if (!theCheck) + { + myRaytraceSceneData.Clear(); + + myIsRaytraceDataValid = Standard_False; + } + else + { + if (myLayersModificationStatus != myView->LayerList().ModificationState()) + { + return UpdateRaytraceGeometry (Standard_False); + } + } + + float* aTransform (NULL); + + // The set of processed structures (reflected to ray-tracing) + // This set is used to remove out-of-date records from the + // hash map of structures + std::set anElements; + + const OpenGl_LayerList& aList = myView->LayerList(); + + for (OpenGl_SequenceOfLayers::Iterator anLayerIt (aList.Layers()); anLayerIt.More(); anLayerIt.Next()) + { + const OpenGl_PriorityList& aPriorityList = anLayerIt.Value(); + + if (aPriorityList.NbStructures() == 0) + continue; + + const OpenGl_ArrayOfStructure& aStructArray = aPriorityList.ArrayOfStructures(); + + for (int anIndex = 0; anIndex < aStructArray.Length(); ++anIndex) + { + OpenGl_SequenceOfStructure::Iterator aStructIt; + + for (aStructIt.Init (aStructArray (anIndex)); aStructIt.More(); aStructIt.Next()) + { + const OpenGl_Structure* aStructure = aStructIt.Value(); + + if (theCheck) + { + if (CheckRaytraceStructure (aStructure)) + { + return UpdateRaytraceGeometry (Standard_False); + } + } + else + { + if (!aStructure->IsRaytracable()) + continue; + + if (aStructure->Transformation()->mat != NULL) + { + if (aTransform == NULL) + aTransform = new float[16]; + + for (int i = 0; i < 4; ++i) + for (int j = 0; j < 4; ++j) + { + aTransform[j * 4 + i] = aStructure->Transformation()->mat[i][j]; + } + } + + AddRaytraceStructure (aStructure, aTransform, anElements); + } + } + } + } + + if (!theCheck) + { + // Actualize the hash map of structures -- remove out-of-date records + std::map::iterator anIter = myStructureStates.begin(); + + while (anIter != myStructureStates.end()) + { + if (anElements.find (anIter->first) == anElements.end()) + { + myStructureStates.erase (anIter++); + } + else + { + ++anIter; + } + } + + // Actualize OpenGL layer list state + myLayersModificationStatus = myView->LayerList().ModificationState(); + + +#ifdef RAY_TRACE_PRINT_INFO + OSD_Timer aTimer; + aTimer.Start(); +#endif + + myBVHBuilder.Build (myRaytraceSceneData); + +#ifdef RAY_TRACE_PRINT_INFO + std::cout << " Build time: " << aTimer.ElapsedTime() << " for " + << myRaytraceSceneData.Triangles.size() / 1000 << "K triangles" << std::endl; +#endif + + const float aScaleFactor = 1.5f; + + myRaytraceSceneRadius = aScaleFactor * + Max ( Max (fabsf (myRaytraceSceneData.AABB.CornerMin().x()), + Max (fabsf (myRaytraceSceneData.AABB.CornerMin().y()), + fabsf (myRaytraceSceneData.AABB.CornerMin().z()))), + Max (fabsf (myRaytraceSceneData.AABB.CornerMax().x()), + Max (fabsf (myRaytraceSceneData.AABB.CornerMax().y()), + fabsf (myRaytraceSceneData.AABB.CornerMax().z()))) ); + + myRaytraceSceneEpsilon = Max (1e-4f, myRaytraceSceneRadius * 1e-4f); + + return WriteRaytraceSceneToDevice(); + } + + delete [] aTransform; + + return Standard_True; +} + +// ======================================================================= +// function : CheckRaytraceStructure +// purpose : Adds OpenGL structure to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::CheckRaytraceStructure (const OpenGl_Structure* theStructure) +{ + if (!theStructure->IsRaytracable()) + { + // Checks to see if all ray-tracable elements were + // removed from the structure + if (theStructure->ModificationState() > 0) + { + theStructure->ResetModificationState(); + return Standard_True; + } + + return Standard_False; + } + + std::map::iterator aStructState = myStructureStates.find (theStructure); + + if (aStructState != myStructureStates.end()) + return aStructState->second != theStructure->ModificationState(); + + return Standard_True; +} + +// ======================================================================= +// function : CreateMaterial +// purpose : Creates ray-tracing material properties +// ======================================================================= +void CreateMaterial (const OPENGL_SURF_PROP& theProp, OpenGl_RaytraceMaterial& theMaterial) +{ + theMaterial.Ambient = OpenGl_RTVec4f (theProp.ambcol.rgb[0] * theProp.amb, + theProp.ambcol.rgb[1] * theProp.amb, + theProp.ambcol.rgb[2] * theProp.amb, + 1.f); + + theMaterial.Diffuse = OpenGl_RTVec4f (theProp.difcol.rgb[0] * theProp.diff, + theProp.difcol.rgb[1] * theProp.diff, + theProp.difcol.rgb[2] * theProp.diff, + 1.f); + + theMaterial.Specular = OpenGl_RTVec4f (theProp.speccol.rgb[0] * theProp.spec, + theProp.speccol.rgb[1] * theProp.spec, + theProp.speccol.rgb[2] * theProp.spec, + theProp.shine); + + theMaterial.Emission = OpenGl_RTVec4f (theProp.emscol.rgb[0] * theProp.emsv, + theProp.emscol.rgb[1] * theProp.emsv, + theProp.emscol.rgb[2] * theProp.emsv, + 1.f); + + // Note: Here we use sub-linear transparency function + // to produce realistic-looking transparency effect + theMaterial.Transparency = OpenGl_RTVec4f (powf (theProp.trans, 0.75f), + 1.f - theProp.trans, + 1.f, + 1.f); + + const float aMaxRefl = Max (theMaterial.Diffuse.x() + theMaterial.Specular.x(), + Max (theMaterial.Diffuse.y() + theMaterial.Specular.y(), + theMaterial.Diffuse.z() + theMaterial.Specular.z())); + + const float aReflectionScale = 0.75f / aMaxRefl; + + theMaterial.Reflection = OpenGl_RTVec4f (theProp.speccol.rgb[0] * theProp.spec, + theProp.speccol.rgb[1] * theProp.spec, + theProp.speccol.rgb[2] * theProp.spec, + 0.f) * aReflectionScale; +} + +// ======================================================================= +// function : AddRaytraceStructure +// purpose : Adds OpenGL structure to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure* theStructure, + const float* theTransform, + std::set& theElements) +{ +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Add Structure" << std::endl; +#endif + + theElements.insert (theStructure); + + if (!theStructure->IsVisible()) + { + myStructureStates[theStructure] = theStructure->ModificationState(); + return Standard_True; + } + + // Get structure material + int aStructMatID = -1; + + if (theStructure->AspectFace() != NULL) + { + aStructMatID = static_cast (myRaytraceSceneData.Materials.size()); + + OpenGl_RaytraceMaterial aStructMaterial; + CreateMaterial (theStructure->AspectFace()->IntFront(), aStructMaterial); + + myRaytraceSceneData.Materials.push_back (aStructMaterial); + } + + OpenGl_ListOfGroup::Iterator anItg (theStructure->Groups()); + + while (anItg.More()) + { + // Get group material + int aGroupMatID = -1; + + if (anItg.Value()->AspectFace() != NULL) + { + aGroupMatID = static_cast (myRaytraceSceneData.Materials.size()); + + OpenGl_RaytraceMaterial aGroupMaterial; + CreateMaterial (anItg.Value()->AspectFace()->IntFront(), aGroupMaterial); + + myRaytraceSceneData.Materials.push_back (aGroupMaterial); + } + + int aMatID = aGroupMatID < 0 ? aStructMatID : aGroupMatID; + + if (aStructMatID < 0 && aGroupMatID < 0) + { + aMatID = static_cast (myRaytraceSceneData.Materials.size()); + + myRaytraceSceneData.Materials.push_back (OpenGl_RaytraceMaterial()); + } + + // Add OpenGL elements from group (only arrays of primitives) + for (const OpenGl_ElementNode* aNode = anItg.Value()->FirstNode(); aNode != NULL; aNode = aNode->next) + { + if (TelNil == aNode->type) + { + OpenGl_AspectFace* anAspect = dynamic_cast (aNode->elem); + + if (anAspect != NULL) + { + aMatID = static_cast (myRaytraceSceneData.Materials.size()); + + OpenGl_RaytraceMaterial aMaterial; + CreateMaterial (anAspect->IntFront(), aMaterial); + + myRaytraceSceneData.Materials.push_back (aMaterial); + } + } + else if (TelParray == aNode->type) + { + OpenGl_PrimitiveArray* aPrimArray = dynamic_cast (aNode->elem); + + if (aPrimArray != NULL) + { + AddRaytracePrimitiveArray (aPrimArray->PArray(), aMatID, theTransform); + } + } + } + + anItg.Next(); + } + + float* aTransform (NULL); + + // Process all connected OpenGL structures + OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures()); + + while (anIts.More()) + { + if (anIts.Value()->Transformation()->mat != NULL) + { + float* aTransform = new float[16]; + + for (int i = 0; i < 4; ++i) + for (int j = 0; j < 4; ++j) + { + aTransform[j * 4 + i] = + anIts.Value()->Transformation()->mat[i][j]; + } + } + + if (anIts.Value()->IsRaytracable()) + AddRaytraceStructure (anIts.Value(), aTransform != NULL ? aTransform : theTransform, theElements); + + anIts.Next(); + } + + delete[] aTransform; + + myStructureStates[theStructure] = theStructure->ModificationState(); + + return Standard_True; +} + +// ======================================================================= +// function : AddRaytracePrimitiveArray +// purpose : Adds OpenGL primitive array to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::AddRaytracePrimitiveArray (const CALL_DEF_PARRAY* theArray, + int theMatID, + const float* theTransform) +{ + if (theArray->type != TelPolygonsArrayType && + theArray->type != TelTrianglesArrayType && + theArray->type != TelQuadranglesArrayType && + theArray->type != TelTriangleFansArrayType && + theArray->type != TelTriangleStripsArrayType && + theArray->type != TelQuadrangleStripsArrayType) + { + return Standard_True; + } + + if (theArray->vertices == NULL) + return Standard_False; + +#ifdef RAY_TRACE_PRINT_INFO + switch (theArray->type) + { + case TelPolygonsArrayType: + std::cout << "\tTelPolygonsArrayType" << std::endl; break; + case TelTrianglesArrayType: + std::cout << "\tTelTrianglesArrayType" << std::endl; break; + case TelQuadranglesArrayType: + std::cout << "\tTelQuadranglesArrayType" << std::endl; break; + case TelTriangleFansArrayType: + std::cout << "\tTelTriangleFansArrayType" << std::endl; break; + case TelTriangleStripsArrayType: + std::cout << "\tTelTriangleStripsArrayType" << std::endl; break; + case TelQuadrangleStripsArrayType: + std::cout << "\tTelQuadrangleStripsArrayType" << std::endl; break; + } +#endif + + // Simple optimization to eliminate possible memory allocations + // during processing of the primitive array vertices + myRaytraceSceneData.Vertices.reserve ( + myRaytraceSceneData.Vertices.size() + theArray->num_vertexs); + + const int aFirstVert = static_cast (myRaytraceSceneData.Vertices.size()); + + for (int aVert = 0; aVert < theArray->num_vertexs; ++aVert) + { + OpenGl_RTVec4f aVertex (theArray->vertices[aVert].xyz[0], + theArray->vertices[aVert].xyz[1], + theArray->vertices[aVert].xyz[2], + 1.f); + + if (theTransform) + aVertex = MatVecMult (theTransform, aVertex); + + myRaytraceSceneData.Vertices.push_back (aVertex); + + myRaytraceSceneData.AABB.Add (aVertex); + } + + myRaytraceSceneData.Normals.reserve ( + myRaytraceSceneData.Normals.size() + theArray->num_vertexs); + + for (int aNorm = 0; aNorm < theArray->num_vertexs; ++aNorm) + { + OpenGl_RTVec4f aNormal; + + // Note: In case of absence of normals, the visualizer + // will use generated geometric normals + + if (theArray->vnormals != NULL) + { + aNormal = OpenGl_RTVec4f (theArray->vnormals[aNorm].xyz[0], + theArray->vnormals[aNorm].xyz[1], + theArray->vnormals[aNorm].xyz[2], + 0.f); + + if (theTransform) + aNormal = MatVecMult (theTransform, aNormal); + } + + myRaytraceSceneData.Normals.push_back (aNormal); + } + + if (theArray->num_bounds > 0) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "\tNumber of bounds = " << theArray->num_bounds << std::endl; +#endif + + int aVertOffset = 0; + + for (int aBound = 0; aBound < theArray->num_bounds; ++aBound) + { + const int aVertNum = theArray->bounds[aBound]; + +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "\tAdd indices from bound " << aBound << ": " << + aVertOffset << ", " << aVertNum << std::endl; +#endif + + if (!AddRaytraceVertexIndices (theArray, aFirstVert, aVertOffset, aVertNum, theMatID)) + { + return Standard_False; + } + + aVertOffset += aVertNum; + } + } + else + { + const int aVertNum = theArray->num_edges > 0 ? theArray->num_edges : theArray->num_vertexs; + +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "\tAdd indices: " << aVertNum << std::endl; +#endif + + return AddRaytraceVertexIndices (theArray, aFirstVert, 0, aVertNum, theMatID); + } + + return Standard_True; +} + +// ======================================================================= +// function : AddRaytraceVertexIndices +// purpose : Adds vertex indices to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::AddRaytraceVertexIndices (const CALL_DEF_PARRAY* theArray, + int theFirstVert, + int theVertOffset, + int theVertNum, + int theMatID) +{ + myRaytraceSceneData.Triangles.reserve (myRaytraceSceneData.Triangles.size() + theVertNum); + + switch (theArray->type) + { + case TelTrianglesArrayType: + return AddRaytraceTriangleArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID); + + case TelQuadranglesArrayType: + return AddRaytraceQuadrangleArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID); + + case TelTriangleFansArrayType: + return AddRaytraceTriangleFanArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID); + + case TelTriangleStripsArrayType: + return AddRaytraceTriangleStripArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID); + + case TelQuadrangleStripsArrayType: + return AddRaytraceQuadrangleStripArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID); + + case TelPolygonsArrayType: + return AddRaytracePolygonArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID); + } + + return Standard_False; +} + +// ======================================================================= +// function : AddRaytraceTriangleArray +// purpose : Adds OpenGL triangle array to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, + int theVertOffset, + int theVertNum, + int theMatID) +{ + if (theVertNum < 3) + return Standard_True; + + if (theArray->num_edges > 0) + { + for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; aVert += 3) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[aVert + 0], + theFirstVert + theArray->edges[aVert + 1], + theFirstVert + theArray->edges[aVert + 2], + theMatID)); + } + } + else + { + for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; aVert += 3) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0, + theFirstVert + aVert + 1, + theFirstVert + aVert + 2, + theMatID)); + } + } + + return Standard_True; +} + +// ======================================================================= +// function : AddRaytraceTriangleFanArray +// purpose : Adds OpenGL triangle fan array to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleFanArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, + int theVertOffset, + int theVertNum, + int theMatID) +{ + if (theVertNum < 3) + return Standard_True; + + if (theArray->num_edges > 0) + { + for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[theVertOffset], + theFirstVert + theArray->edges[aVert + 1], + theFirstVert + theArray->edges[aVert + 2], + theMatID)); + } + } + else + { + for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset, + theFirstVert + aVert + 1, + theFirstVert + aVert + 2, + theMatID)); + } + } + + return Standard_True; +} + +// ======================================================================= +// function : AddRaytraceTriangleStripArray +// purpose : Adds OpenGL triangle strip array to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleStripArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, + int theVertOffset, + int theVertNum, + int theMatID) +{ + if (theVertNum < 3) + return Standard_True; + + if (theArray->num_edges > 0) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i ( + theFirstVert + theArray->edges[theVertOffset + 0], + theFirstVert + theArray->edges[theVertOffset + 1], + theFirstVert + theArray->edges[theVertOffset + 2], + theMatID)); + + for (int aVert = theVertOffset + 1, aTriNum = 1; aVert < theVertOffset + theVertNum - 2; ++aVert, ++aTriNum) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i ( + theFirstVert + theArray->edges[aVert + (aTriNum % 2) ? 1 : 0], + theFirstVert + theArray->edges[aVert + (aTriNum % 2) ? 0 : 1], + theFirstVert + theArray->edges[aVert + 2], + theMatID)); + } + } + else + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset + 0, + theFirstVert + theVertOffset + 1, + theFirstVert + theVertOffset + 2, + theMatID)); + + for (int aVert = theVertOffset + 1, aTriNum = 1; aVert < theVertOffset + theVertNum - 2; ++aVert, ++aTriNum) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + ( aTriNum % 2 ) ? 1 : 0, + theFirstVert + aVert + ( aTriNum % 2 ) ? 0 : 1, + theFirstVert + aVert + 2, + theMatID)); + } + } + + return Standard_True; +} + +// ======================================================================= +// function : AddRaytraceQuadrangleArray +// purpose : Adds OpenGL quad array to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, + int theVertOffset, + int theVertNum, + int theMatID) +{ + if (theVertNum < 4) + return Standard_True; + + if (theArray->num_edges > 0) + { + for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 3; aVert += 4) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[aVert + 0], + theFirstVert + theArray->edges[aVert + 1], + theFirstVert + theArray->edges[aVert + 2], + theMatID)); + + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[aVert + 0], + theFirstVert + theArray->edges[aVert + 2], + theFirstVert + theArray->edges[aVert + 3], + theMatID)); + } + } + else + { + for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 3; aVert += 4) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0, + theFirstVert + aVert + 1, + theFirstVert + aVert + 2, + theMatID)); + + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0, + theFirstVert + aVert + 2, + theFirstVert + aVert + 3, + theMatID)); + } + } + + return Standard_True; +} + +// ======================================================================= +// function : AddRaytraceQuadrangleStripArray +// purpose : Adds OpenGL quad strip array to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleStripArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, + int theVertOffset, + int theVertNum, + int theMatID) +{ + if (theVertNum < 4) + return Standard_True; + + if (theArray->num_edges > 0) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i ( + theFirstVert + theArray->edges[theVertOffset + 0], + theFirstVert + theArray->edges[theVertOffset + 1], + theFirstVert + theArray->edges[theVertOffset + 2], + theMatID)); + + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i ( + theFirstVert + theArray->edges[theVertOffset + 1], + theFirstVert + theArray->edges[theVertOffset + 3], + theFirstVert + theArray->edges[theVertOffset + 2], + theMatID)); + + for (int aVert = theVertOffset + 2; aVert < theVertOffset + theVertNum - 3; aVert += 2) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i ( + theFirstVert + theArray->edges[aVert + 0], + theFirstVert + theArray->edges[aVert + 1], + theFirstVert + theArray->edges[aVert + 2], + theMatID)); + + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i ( + theFirstVert + theArray->edges[aVert + 1], + theFirstVert + theArray->edges[aVert + 3], + theFirstVert + theArray->edges[aVert + 2], + theMatID)); + } + } + else + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + 0, + theFirstVert + 1, + theFirstVert + 2, + theMatID)); + + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + 1, + theFirstVert + 3, + theFirstVert + 2, + theMatID)); + + for (int aVert = theVertOffset + 2; aVert < theVertOffset + theVertNum - 3; aVert += 2) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0, + theFirstVert + aVert + 1, + theFirstVert + aVert + 2, + theMatID)); + + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 1, + theFirstVert + aVert + 3, + theFirstVert + aVert + 2, + theMatID)); + } + } + + return Standard_True; +} + +// ======================================================================= +// function : AddRaytracePolygonArray +// purpose : Adds OpenGL polygon array to ray-traced scene geometry +// ======================================================================= +Standard_Boolean OpenGl_Workspace::AddRaytracePolygonArray (const CALL_DEF_PARRAY* theArray, + int theFirstVert, + int theVertOffset, + int theVertNum, + int theMatID) +{ + if (theArray->num_vertexs < 3) + return Standard_True; + + if (theArray->edges != NULL) + { + for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[theVertOffset], + theFirstVert + theArray->edges[aVert + 1], + theFirstVert + theArray->edges[aVert + 2], + theMatID)); + } + } + else + { + for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert) + { + myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset, + theFirstVert + aVert + 1, + theFirstVert + aVert + 2, + theMatID)); + } + } + + return Standard_True; +} + +// ======================================================================= +// function : UpdateRaytraceLightSources +// purpose : Updates 3D scene light sources for ray-tracing +// ======================================================================= +Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble theInvModelView[16]) +{ + myRaytraceSceneData.LightSources.clear(); + + OpenGl_ListOfLight::Iterator anItl (myView->LightList()); + + OpenGl_RTVec4f aAmbient (0.f, 0.f, 0.f, 0.f); + + for (; anItl.More(); anItl.Next()) + { + const OpenGl_Light &aLight = anItl.Value(); + + if (aLight.type == TLightAmbient) + { + aAmbient += OpenGl_RTVec4f (aLight.col.rgb[0], + aLight.col.rgb[1], + aLight.col.rgb[2], + 0.f); + continue; + } + + OpenGl_RTVec4f aDiffuse (aLight.col.rgb[0], + aLight.col.rgb[1], + aLight.col.rgb[2], + 1.f); + + OpenGl_RTVec4f aPosition (-aLight.dir[0], + -aLight.dir[1], + -aLight.dir[2], + 0.f); + + if (aLight.type != TLightDirectional) + { + aPosition = OpenGl_RTVec4f (aLight.pos[0], + aLight.pos[1], + aLight.pos[2], + 1.f); + } + + if (aLight.HeadLight) + { + aPosition = MatVecMult (theInvModelView, aPosition); + } + + myRaytraceSceneData.LightSources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition)); + } + + if (myRaytraceSceneData.LightSources.size() > 0) + { + myRaytraceSceneData.LightSources.front().Ambient += aAmbient; + } + else + { + myRaytraceSceneData.LightSources.push_back (OpenGl_RaytraceLight (OpenGl_RTVec4f (aAmbient.xyz(), -1.0f))); + } + + cl_int anError = CL_SUCCESS; + + if (myRaytraceLightSourceBuffer != NULL) + clReleaseMemObject (myRaytraceLightSourceBuffer); + + const size_t myLightBufferSize = myRaytraceSceneData.LightSources.size() > 0 + ? myRaytraceSceneData.LightSources.size() + : 1; + + myRaytraceLightSourceBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, + myLightBufferSize * sizeof(OpenGl_RaytraceLight), + NULL, &anError); + + if (myRaytraceSceneData.LightSources.size() > 0) + { + const void* aDataPtr = myRaytraceSceneData.LightSources.front().Packed(); + anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0, + myLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr, + 0, NULL, NULL); + } + +#ifdef RAY_TRACE_PRINT_INFO + if (anError != CL_SUCCESS) + { + std::cout << "Error! Failed to set light sources!"; + + return Standard_False; + } +#endif + + return Standard_True; +} + +// ======================================================================= +// function : CheckOpenCL +// purpose : Checks OpenCL dynamic library availability +// ======================================================================= +Standard_Boolean CheckOpenCL() +{ +#if defined ( _WIN32 ) + + __try + { + cl_uint aNbPlatforms; + clGetPlatformIDs (0, NULL, &aNbPlatforms); + } + __except (EXCEPTION_EXECUTE_HANDLER) + { + return Standard_False; + } + +#endif + + return Standard_True; +} + +// ======================================================================= +// function : InitOpenCL +// purpose : Initializes OpenCL objects +// ======================================================================= +Standard_Boolean OpenGl_Workspace::InitOpenCL() +{ + if (myComputeInitStatus != OpenGl_CLIS_NONE) + { + return myComputeInitStatus == OpenGl_CLIS_INIT; + } + + if (!CheckOpenCL()) + { + myComputeInitStatus = OpenGl_CLIS_FAIL; // fail to load OpenCL library + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + "Failed to load OpenCL dynamic library!"); + return Standard_False; + } + + // Obtain the list of platforms available + cl_uint aNbPlatforms = 0; + cl_int anError = clGetPlatformIDs (0, NULL, &aNbPlatforms); + cl_platform_id* aPlatforms = (cl_platform_id* )alloca (aNbPlatforms * sizeof(cl_platform_id)); + anError |= clGetPlatformIDs (aNbPlatforms, aPlatforms, NULL); + if (anError != CL_SUCCESS + || aNbPlatforms == 0) + { + myComputeInitStatus = OpenGl_CLIS_FAIL; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + "No any OpenCL platform installed!"); + return Standard_False; + } + + // Note: We try to find NVIDIA or AMD platforms with GPU devices! + cl_platform_id aPrefPlatform = NULL; + for (cl_uint aPlatIter = 0; aPlatIter < aNbPlatforms; ++aPlatIter) + { + char aName[256]; + anError = clGetPlatformInfo (aPlatforms[aPlatIter], CL_PLATFORM_NAME, + sizeof(aName), aName, NULL); + if (anError != CL_SUCCESS) + { + continue; + } + + if (strncmp (aName, "NVIDIA", strlen ("NVIDIA")) == 0) + { + aPrefPlatform = aPlatforms[aPlatIter]; + + // Use optimizations for NVIDIA GPUs + myIsAmdComputePlatform = Standard_False; + } + else if (strncmp (aName, "AMD", strlen ("AMD")) == 0) + { + aPrefPlatform = (aPrefPlatform == NULL) + ? aPlatforms[aPlatIter] + : aPrefPlatform; + + // Use optimizations for ATI/AMD platform + myIsAmdComputePlatform = Standard_True; + } + } + + if (aPrefPlatform == NULL) + { + aPrefPlatform = aPlatforms[0]; + } + + // Obtain the list of devices available in the selected platform + cl_uint aNbDevices = 0; + anError = clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU, + 0, NULL, &aNbDevices); + + cl_device_id* aDevices = (cl_device_id* )alloca (aNbDevices * sizeof(cl_device_id)); + anError |= clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU, + aNbDevices, aDevices, NULL); + if (anError != CL_SUCCESS) + { + myComputeInitStatus = OpenGl_CLIS_FAIL; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + "Failed to get OpenCL GPU device!"); + return Standard_False; + } + + // Note: Simply get first available GPU + cl_device_id aDevice = aDevices[0]; + + // detect old contexts + char aVerClStr[256]; + clGetDeviceInfo (aDevice, CL_DEVICE_VERSION, + sizeof(aVerClStr), aVerClStr, NULL); + aVerClStr[strlen ("OpenCL 1.0")] = '\0'; + const bool isVer10 = strncmp (aVerClStr, "OpenCL 1.0", strlen ("OpenCL 1.0")) == 0; + + // Create OpenCL context + cl_context_properties aCtxProp[] = + { + #if defined(__APPLE__) && !defined(MACOSX_USE_GLX) + CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, + (cl_context_properties )CGLGetShareGroup (CGLGetCurrentContext()), + #elif defined(_WIN32) + CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform, + CL_GL_CONTEXT_KHR, (cl_context_properties )wglGetCurrentContext(), + CL_WGL_HDC_KHR, (cl_context_properties )wglGetCurrentDC(), + #else + CL_GL_CONTEXT_KHR, (cl_context_properties )glXGetCurrentContext(), + CL_GLX_DISPLAY_KHR, (cl_context_properties )glXGetCurrentDisplay(), + CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform, + #endif + 0 + }; + + myComputeContext = clCreateContext (aCtxProp, + #if defined(__APPLE__) && !defined(MACOSX_USE_GLX) + 0, NULL, // device will be taken from GL context + #else + 1, &aDevice, + #endif + NULL, NULL, &anError); + if (anError != CL_SUCCESS) + { + myComputeInitStatus = OpenGl_CLIS_FAIL; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + "Failed to initialize OpenCL context!"); + return Standard_False; + } + + // Create OpenCL program + const char* aSources[] = + { + isVer10 ? "#define M_PI_F ( float )( 3.14159265359f )\n" : "", + THE_RAY_TRACE_OPENCL_SOURCE + }; + myRaytraceProgram = clCreateProgramWithSource (myComputeContext, 2, + aSources, NULL, &anError); + if (anError != CL_SUCCESS) + { + myComputeInitStatus = OpenGl_CLIS_FAIL; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + "Failed to create OpenCL ray-tracing program!"); + return Standard_False; + } + + anError = clBuildProgram (myRaytraceProgram, 0, + NULL, NULL, NULL, NULL); + { + // Fetch build log + size_t aLogLen = 0; + cl_int aResult = clGetProgramBuildInfo (myRaytraceProgram, aDevice, + CL_PROGRAM_BUILD_LOG, 0, NULL, &aLogLen); + + char* aBuildLog = (char* )alloca (aLogLen); + aResult |= clGetProgramBuildInfo (myRaytraceProgram, aDevice, + CL_PROGRAM_BUILD_LOG, aLogLen, aBuildLog, NULL); + if (aResult == CL_SUCCESS) + { + if (anError != CL_SUCCESS) + { + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + aBuildLog); + } + else + { + #ifdef RAY_TRACE_PRINT_INFO + std::cout << aBuildLog << std::endl; + #endif + } + } + } + + if (anError != CL_SUCCESS) + { + return Standard_False; + } + + // Create OpenCL ray tracing kernels + myRaytraceRenderKernel = clCreateKernel (myRaytraceProgram, "Main", &anError); + if (anError != CL_SUCCESS) + { + myComputeInitStatus = OpenGl_CLIS_FAIL; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + "Failed to create OpenCL ray-tracing kernel!"); + return Standard_False; + } + + myRaytraceSmoothKernel = clCreateKernel (myRaytraceProgram, "MainAntialiased", &anError); + if (anError != CL_SUCCESS) + { + myComputeInitStatus = OpenGl_CLIS_FAIL; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + "Failed to create OpenCL ray-tracing kernel!"); + return Standard_False; + } + + // Create OpenCL command queue + // Note: For profiling set CL_QUEUE_PROFILING_ENABLE + cl_command_queue_properties aProps = CL_QUEUE_PROFILING_ENABLE; + + myRaytraceQueue = clCreateCommandQueue (myComputeContext, aDevice, aProps, &anError); + if (anError != CL_SUCCESS) + { + myComputeInitStatus = OpenGl_CLIS_FAIL; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + "Failed to create OpenCL command queue!"); + + return Standard_False; + } + + myComputeInitStatus = OpenGl_CLIS_INIT; // initialized in normal way + return Standard_True; +} + +// ======================================================================= +// function : GetOpenClDeviceInfo +// purpose : Returns information about device used for computations +// ======================================================================= +Standard_Boolean OpenGl_Workspace::GetOpenClDeviceInfo (NCollection_DataMap& theInfo) const +{ + theInfo.Clear(); + if (myComputeContext == NULL) + { + return Standard_False; + } + + size_t aDevicesSize = 0; + cl_int anError = clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, 0, NULL, &aDevicesSize); + cl_device_id* aDevices = (cl_device_id* )alloca (aDevicesSize); + anError |= clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, aDevicesSize, aDevices, NULL); + if (anError != CL_SUCCESS) + { + return Standard_False; + } + + char aDeviceName[256]; + anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_NAME, sizeof(aDeviceName), aDeviceName, NULL); + theInfo.Bind ("Name", aDeviceName); + + char aDeviceVendor[256]; + anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_VENDOR, sizeof(aDeviceVendor), aDeviceVendor, NULL); + theInfo.Bind ("Vendor", aDeviceVendor); + + cl_device_type aDeviceType; + anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_TYPE, sizeof(aDeviceType), &aDeviceType, NULL); + theInfo.Bind ("Type", aDeviceType == CL_DEVICE_TYPE_GPU ? "GPU" : "CPU"); + return Standard_True; +} + +// ======================================================================= +// function : ReleaseOpenCL +// purpose : Releases resources of OpenCL objects +// ======================================================================= +void OpenGl_Workspace::ReleaseOpenCL() +{ + clReleaseKernel (myRaytraceRenderKernel); + clReleaseKernel (myRaytraceSmoothKernel); + + clReleaseProgram (myRaytraceProgram); + clReleaseCommandQueue (myRaytraceQueue); + + clReleaseMemObject (myRaytraceOutputImage); + clReleaseMemObject (myRaytraceEnvironment); + clReleaseMemObject (myRaytraceOutputImageSmooth); + + clReleaseMemObject (myRaytraceVertexBuffer); + clReleaseMemObject (myRaytraceNormalBuffer); + clReleaseMemObject (myRaytraceTriangleBuffer); + + clReleaseMemObject (myRaytraceMaterialBuffer); + clReleaseMemObject (myRaytraceLightSourceBuffer); + + clReleaseMemObject (myRaytraceNodeMinPointBuffer); + clReleaseMemObject (myRaytraceNodeMaxPointBuffer); + clReleaseMemObject (myRaytraceNodeDataRcrdBuffer); + + clReleaseContext (myComputeContext); + + if (glIsTexture (*myRaytraceOutputTexture)) + glDeleteTextures (2, myRaytraceOutputTexture); +} + +// ======================================================================= +// function : ResizeRaytraceOutputBuffer +// purpose : Resizes OpenCL output image +// ======================================================================= +Standard_Boolean OpenGl_Workspace::ResizeRaytraceOutputBuffer (const cl_int theSizeX, + const cl_int theSizeY) +{ + if (myComputeContext == NULL) + { + return Standard_False; + } + + bool toResize = true; + GLint aSizeX = -1; + GLint aSizeY = -1; + if (*myRaytraceOutputTexture != 0) + { + if (!myGlContext->IsGlGreaterEqual (2, 1)) + { + return Standard_False; + } + + glBindTexture (GL_TEXTURE_RECTANGLE, *myRaytraceOutputTexture); + + glGetTexLevelParameteriv (GL_TEXTURE_RECTANGLE, 0, GL_TEXTURE_WIDTH, &aSizeX); + glGetTexLevelParameteriv (GL_TEXTURE_RECTANGLE, 0, GL_TEXTURE_HEIGHT, &aSizeY); + + toResize = (aSizeX != theSizeX) || (aSizeY != theSizeY); + if (toResize) + { + glDeleteTextures (2, myRaytraceOutputTexture); + } + } + if (!toResize) + { + return Standard_True; + } + + glGenTextures (2, myRaytraceOutputTexture); + for (int aTexIter = 0; aTexIter < 2; ++aTexIter) + { + glBindTexture (GL_TEXTURE_RECTANGLE, myRaytraceOutputTexture[aTexIter]); + + glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_WRAP_S, GL_CLAMP); + glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_WRAP_T, GL_CLAMP); + glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_WRAP_R, GL_CLAMP); + + glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + + glTexImage2D (GL_TEXTURE_RECTANGLE, 0, GL_RGBA32F, + theSizeX, theSizeY, 0, + GL_RGBA, GL_FLOAT, NULL); + } + + cl_int anError = CL_SUCCESS; + + if (myRaytraceOutputImage != NULL) + { + clReleaseMemObject (myRaytraceOutputImage); + } + if (myRaytraceOutputImageSmooth != NULL) + { + clReleaseMemObject (myRaytraceOutputImageSmooth); + } + + myRaytraceOutputImage = clCreateFromGLTexture2D (myComputeContext, CL_MEM_READ_WRITE, + GL_TEXTURE_RECTANGLE, 0, + myRaytraceOutputTexture[0], &anError); + if (anError != CL_SUCCESS) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error! Failed to create output image!" << std::endl; +#endif + return Standard_False; + } + + myRaytraceOutputImageSmooth = clCreateFromGLTexture2D (myComputeContext, CL_MEM_READ_WRITE, + GL_TEXTURE_RECTANGLE, 0, + myRaytraceOutputTexture[1], &anError); + if (anError != CL_SUCCESS) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error! Failed to create anti-aliased output image!" << std::endl; +#endif + return Standard_False; + } + + return Standard_True; +} + +// ======================================================================= +// function : WriteRaytraceSceneToDevice +// purpose : Writes scene geometry to OpenCl device +// ======================================================================= +Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice() +{ + if (myComputeContext == NULL) + return Standard_False; + + cl_int anError = CL_SUCCESS; + + if (myRaytraceNormalBuffer != NULL) + anError |= clReleaseMemObject (myRaytraceNormalBuffer); + + if (myRaytraceVertexBuffer != NULL) + anError |= clReleaseMemObject (myRaytraceVertexBuffer); + + if (myRaytraceTriangleBuffer != NULL) + anError |= clReleaseMemObject (myRaytraceTriangleBuffer); + + if (myRaytraceNodeMinPointBuffer != NULL) + anError |= clReleaseMemObject (myRaytraceNodeMinPointBuffer); + + if (myRaytraceNodeMaxPointBuffer != NULL) + anError |= clReleaseMemObject (myRaytraceNodeMaxPointBuffer); + + if (myRaytraceNodeDataRcrdBuffer != NULL) + anError |= clReleaseMemObject (myRaytraceNodeDataRcrdBuffer); + + if (myRaytraceMaterialBuffer != NULL) + anError |= clReleaseMemObject (myRaytraceMaterialBuffer); + + if (anError != CL_SUCCESS) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error! Failed to release OpenCL scene buffers!" << std::endl; +#endif + return Standard_False; + } + + // Create geometry buffers + cl_int anErrorTemp = CL_SUCCESS; + const size_t myVertexBufferSize = myRaytraceSceneData.Vertices.size() > 0 + ? myRaytraceSceneData.Vertices.size() : 1; + + myRaytraceVertexBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, + myVertexBufferSize * sizeof(cl_float4), NULL, &anErrorTemp); + anError |= anErrorTemp; + + const size_t myNormalBufferSize = myRaytraceSceneData.Normals.size() > 0 + ? myRaytraceSceneData.Normals.size() : 1; + myRaytraceNormalBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, + myNormalBufferSize * sizeof(cl_float4), NULL, &anErrorTemp); + anError |= anErrorTemp; + + const size_t myTriangleBufferSize = myRaytraceSceneData.Triangles.size() > 0 + ? myRaytraceSceneData.Triangles.size() : 1; + myRaytraceTriangleBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, + myTriangleBufferSize * sizeof(cl_int4), NULL, &anErrorTemp); + anError |= anErrorTemp; + if (anError != CL_SUCCESS) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error! Failed to create OpenCL geometry buffers!" << std::endl; +#endif + return Standard_False; + } + + // Create material buffer + const size_t myMaterialBufferSize = myRaytraceSceneData.Materials.size() > 0 + ? myRaytraceSceneData.Materials.size() : 1; + myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, + myMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL, + &anErrorTemp); + if (anErrorTemp != CL_SUCCESS) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error! Failed to create OpenCL material buffer!" << std::endl; +#endif + return Standard_False; + } + + // Create BVH buffers + OpenGl_BVH aTree = myBVHBuilder.Tree(); + const size_t myNodeMinPointBufferSize = aTree.MinPointBuffer().size() > 0 + ? aTree.MinPointBuffer().size() : 1; + myRaytraceNodeMinPointBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, + myNodeMinPointBufferSize * sizeof(cl_float4), NULL, + &anErrorTemp); + anError |= anErrorTemp; + + const size_t myNodeMaxPointBufferSize = aTree.MaxPointBuffer().size() > 0 + ? aTree.MaxPointBuffer().size() : 1; + myRaytraceNodeMaxPointBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, + myNodeMaxPointBufferSize * sizeof(cl_float4), NULL, + &anError); + anError |= anErrorTemp; + + const size_t myNodeDataRecordBufferSize = aTree.DataRcrdBuffer().size() > 0 + ? aTree.DataRcrdBuffer().size() : 1; + myRaytraceNodeDataRcrdBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, + myNodeDataRecordBufferSize * sizeof(cl_int4), NULL, + &anError); + anError |= anErrorTemp; + if (anError != CL_SUCCESS) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error! Failed to create OpenCL BVH buffers!" << std::endl; +#endif + return Standard_False; + } + + // Write scene geometry buffers + if (myRaytraceSceneData.Triangles.size() > 0) + { + anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceVertexBuffer, CL_FALSE, + 0, myRaytraceSceneData.Vertices.size() * sizeof(cl_float4), + &myRaytraceSceneData.Vertices.front(), + 0, NULL, NULL); + anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNormalBuffer, CL_FALSE, + 0, myRaytraceSceneData.Normals.size() * sizeof(cl_float4), + &myRaytraceSceneData.Normals.front(), + 0, NULL, NULL); + anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceTriangleBuffer, CL_FALSE, + 0, myRaytraceSceneData.Triangles.size() * sizeof(cl_int4), + &myRaytraceSceneData.Triangles.front(), + 0, NULL, NULL); + if (anError != CL_SUCCESS) + { + #ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error! Failed to write OpenCL geometry buffers!" << std::endl; + #endif + return Standard_False; + } + } + + // Write BVH buffers + if (aTree.DataRcrdBuffer().size() > 0) + { + anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeMinPointBuffer, CL_FALSE, + 0, aTree.MinPointBuffer().size() * sizeof(cl_float4), + &aTree.MinPointBuffer().front(), + 0, NULL, NULL); + anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeMaxPointBuffer, CL_FALSE, + 0, aTree.MaxPointBuffer().size() * sizeof(cl_float4), + &aTree.MaxPointBuffer().front(), + 0, NULL, NULL); + anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeDataRcrdBuffer, CL_FALSE, + 0, aTree.DataRcrdBuffer().size() * sizeof(cl_int4), + &aTree.DataRcrdBuffer().front(), + 0, NULL, NULL); + if (anError != CL_SUCCESS) + { + #ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error! Failed to write OpenCL BVH buffers!" << std::endl; + #endif + return Standard_False; + } + } + + // Write material buffers + if (myRaytraceSceneData.Materials.size() > 0) + { + const size_t aSize = myRaytraceSceneData.Materials.size(); + const void* aDataPtr = myRaytraceSceneData.Materials.front().Packed(); + + anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceMaterialBuffer, CL_FALSE, + 0, aSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr, + 0, NULL, NULL); + if (anError != CL_SUCCESS) + { + #ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error! Failed to write OpenCL material buffer!" << std::endl; + #endif + return Standard_False; + } + } + + anError |= clFinish (myRaytraceQueue); +#ifdef RAY_TRACE_PRINT_INFO + if (anError != CL_SUCCESS) + std::cout << "Error! Failed to set scene data buffers!" << std::endl; +#endif + + if (anError == CL_SUCCESS) + myIsRaytraceDataValid = myRaytraceSceneData.Triangles.size() > 0; + +#ifdef RAY_TRACE_PRINT_INFO + + float aMemUsed = static_cast ( + myRaytraceSceneData.Materials.size()) * sizeof (OpenGl_RaytraceMaterial); + + aMemUsed += static_cast ( + myRaytraceSceneData.Triangles.size() * sizeof (OpenGl_RTVec4i) + + myRaytraceSceneData.Vertices.size() * sizeof (OpenGl_RTVec4f) + + myRaytraceSceneData.Normals.size() * sizeof (OpenGl_RTVec4f)); + + aMemUsed += static_cast ( + aTree.MinPointBuffer().size() * sizeof (OpenGl_RTVec4f) + + aTree.MaxPointBuffer().size() * sizeof (OpenGl_RTVec4f) + + aTree.DataRcrdBuffer().size() * sizeof (OpenGl_RTVec4i)); + + std::cout << "GPU memory used (Mb): " << aMemUsed / 1e6f << std::endl; + +#endif + + myRaytraceSceneData.Clear(); + + myBVHBuilder.CleanUp(); + + return (CL_SUCCESS == anError); +} + +#define OPENCL_GROUP_SIZE_TEST_ + +// ======================================================================= +// function : RunRaytraceOpenCLKernels +// purpose : Runs OpenCL ray-tracing kernels +// ======================================================================= +Standard_Boolean OpenGl_Workspace::RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView, + const GLfloat theOrigins[16], + const GLfloat theDirects[16], + const int theSizeX, + const int theSizeY) +{ + if (myRaytraceRenderKernel == NULL || myRaytraceQueue == NULL) + return Standard_False; + + //////////////////////////////////////////////////////////// + // Set kernel arguments + + cl_uint anIndex = 0; + cl_int anError = 0; + + anError = clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceOutputImage); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceEnvironment); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceNodeMinPointBuffer); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceNodeMaxPointBuffer); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceNodeDataRcrdBuffer); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceLightSourceBuffer); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceMaterialBuffer); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceVertexBuffer); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceNormalBuffer); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_mem), &myRaytraceTriangleBuffer); + + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_float16), theOrigins); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_float16), theDirects); + + cl_int aLightCount = static_cast (myRaytraceSceneData.LightSources.size()); + + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_int), &aLightCount); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_float), &myRaytraceSceneEpsilon); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_float), &myRaytraceSceneRadius); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_int), &theCView.IsShadowsEnabled); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_int), &theCView.IsReflectionsEnabled); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_int), &theSizeX); + anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++, + sizeof(cl_int), &theSizeY); + if (anError != CL_SUCCESS) + { + const TCollection_ExtendedString aMsg = "Error! Failed to set arguments of ray-tracing kernel!"; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + aMsg); + return Standard_False; + } + + // Note: second-pass 'smoothing' kernel runs only if anti-aliasing is enabled + if (theCView.IsAntialiasingEnabled) + { + anIndex = 0; + anError = clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceOutputImage); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceOutputImageSmooth); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceEnvironment); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceNodeMinPointBuffer); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceNodeMaxPointBuffer); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceNodeDataRcrdBuffer); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceLightSourceBuffer); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceMaterialBuffer); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceVertexBuffer); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceNormalBuffer); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_mem), &myRaytraceTriangleBuffer); + + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_float16), theOrigins); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_float16), theDirects); + + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_int), &aLightCount); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_float), &myRaytraceSceneEpsilon); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_float), &myRaytraceSceneRadius); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_int), &theCView.IsShadowsEnabled); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_int), &theCView.IsReflectionsEnabled); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_int), &theSizeX); + anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++, + sizeof(cl_int), &theSizeY); + if (anError != CL_SUCCESS) + { + const TCollection_ExtendedString aMsg = "Error! Failed to set arguments of 'smoothing' kernel!"; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + aMsg); + return Standard_False; + } + } + + // Set work size + size_t aLocSizeRender[] = { myIsAmdComputePlatform ? 2 : 4, 32 }; + +#ifdef OPENCL_GROUP_SIZE_TEST + for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1 ) + for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1 ) +#endif + { +#ifdef OPENCL_GROUP_SIZE_TEST + aLocSizeRender[0] = aLocX; + aLocSizeRender[1] = aLocY; +#endif + + size_t aWorkSizeX = theSizeX; + if (aWorkSizeX % aLocSizeRender[0] != 0) + aWorkSizeX += aLocSizeRender[0] - aWorkSizeX % aLocSizeRender[0]; + + size_t aWokrSizeY = theSizeY; + if (aWokrSizeY % aLocSizeRender[1] != 0 ) + aWokrSizeY += aLocSizeRender[1] - aWokrSizeY % aLocSizeRender[1]; + + size_t aGlbSizeRender[] = { aWorkSizeX, aWokrSizeY }; + + // Run kernel + cl_event anEvent (NULL), anEventSmooth (NULL); + anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceRenderKernel, + 2, NULL, aGlbSizeRender, aLocSizeRender, + 0, NULL, &anEvent); + if (anError != CL_SUCCESS) + { + const TCollection_ExtendedString aMsg = "Error! Failed to execute the ray-tracing kernel!"; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + aMsg); + return Standard_False; + } + clWaitForEvents (1, &anEvent); + + if (theCView.IsAntialiasingEnabled) + { + size_t aLocSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4, + myIsAmdComputePlatform ? 8 : 32 }; + +#ifdef OPENCL_GROUP_SIZE_TEST + aLocSizeSmooth[0] = aLocX; + aLocSizeSmooth[1] = aLocY; +#endif + + aWorkSizeX = theSizeX; + if (aWorkSizeX % aLocSizeSmooth[0] != 0) + aWorkSizeX += aLocSizeSmooth[0] - aWorkSizeX % aLocSizeSmooth[0]; + + size_t aWokrSizeY = theSizeY; + if (aWokrSizeY % aLocSizeSmooth[1] != 0 ) + aWokrSizeY += aLocSizeSmooth[1] - aWokrSizeY % aLocSizeSmooth[1]; + + size_t aGlbSizeSmooth [] = { aWorkSizeX, aWokrSizeY }; + anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceSmoothKernel, + 2, NULL, aGlbSizeSmooth, aLocSizeSmooth, + 0, NULL, &anEventSmooth); + clWaitForEvents (1, &anEventSmooth); + + if (anError != CL_SUCCESS) + { + const TCollection_ExtendedString aMsg = "Error! Failed to execute the 'smoothing' kernel!"; + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, + 0, + GL_DEBUG_SEVERITY_HIGH_ARB, + aMsg); + return Standard_False; + } + } + + // Get the profiling data +#if defined (RAY_TRACE_PRINT_INFO) || defined(OPENCL_GROUP_SIZE_TEST) + + cl_ulong aTimeStart, + aTimeFinal; + + clGetEventProfilingInfo (anEvent, CL_PROFILING_COMMAND_START, + sizeof(aTimeStart), &aTimeStart, NULL); + clGetEventProfilingInfo (anEvent, CL_PROFILING_COMMAND_END, + sizeof(aTimeFinal), &aTimeFinal, NULL); + std::cout << "\tRender time (ms): " << ( aTimeFinal - aTimeStart ) / 1e6f << std::endl; + + if (theCView.IsAntialiasingEnabled) + { + clGetEventProfilingInfo (anEventSmooth, CL_PROFILING_COMMAND_START, + sizeof(aTimeStart), &aTimeStart, NULL); + clGetEventProfilingInfo (anEventSmooth, CL_PROFILING_COMMAND_END, + sizeof(aTimeFinal), &aTimeFinal, NULL); + std::cout << "\tSmoothing time (ms): " << ( aTimeFinal - aTimeStart ) / 1e6f << std::endl; + } +#endif + + if (anEvent != NULL) + clReleaseEvent (anEvent); + + if (anEventSmooth != NULL) + clReleaseEvent (anEventSmooth); + } + + return Standard_True; +} + +// ======================================================================= +// function : ComputeInverseMatrix +// purpose : Computes inversion of 4x4 floating-point matrix +// ======================================================================= +template +void ComputeInverseMatrix (const T m[16], T inv[16]) +{ + inv[ 0] = m[ 5] * (m[10] * m[15] - m[11] * m[14]) - + m[ 9] * (m[ 6] * m[15] - m[ 7] * m[14]) - + m[13] * (m[ 7] * m[10] - m[ 6] * m[11]); + + inv[ 1] = m[ 1] * (m[11] * m[14] - m[10] * m[15]) - + m[ 9] * (m[ 3] * m[14] - m[ 2] * m[15]) - + m[13] * (m[ 2] * m[11] - m[ 3] * m[10]); + + inv[ 2] = m[ 1] * (m[ 6] * m[15] - m[ 7] * m[14]) - + m[ 5] * (m[ 2] * m[15] - m[ 3] * m[14]) - + m[13] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]); + + inv[ 3] = m[ 1] * (m[ 7] * m[10] - m[ 6] * m[11]) - + m[ 5] * (m[ 3] * m[10] - m[ 2] * m[11]) - + m[ 9] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]); + + inv[ 4] = m[ 4] * (m[11] * m[14] - m[10] * m[15]) - + m[ 8] * (m[ 7] * m[14] - m[ 6] * m[15]) - + m[12] * (m[ 6] * m[11] - m[ 7] * m[10]); + + inv[ 5] = m[ 0] * (m[10] * m[15] - m[11] * m[14]) - + m[ 8] * (m[ 2] * m[15] - m[ 3] * m[14]) - + m[12] * (m[ 3] * m[10] - m[ 2] * m[11]); + + inv[ 6] = m[ 0] * (m[ 7] * m[14] - m[ 6] * m[15]) - + m[ 4] * (m[ 3] * m[14] - m[ 2] * m[15]) - + m[12] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]); + + inv[ 7] = m[ 0] * (m[ 6] * m[11] - m[ 7] * m[10]) - + m[ 4] * (m[ 2] * m[11] - m[ 3] * m[10]) - + m[ 8] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]); + + inv[ 8] = m[ 4] * (m[ 9] * m[15] - m[11] * m[13]) - + m[ 8] * (m[ 5] * m[15] - m[ 7] * m[13]) - + m[12] * (m[ 7] * m[ 9] - m[ 5] * m[11]); + + inv[ 9] = m[ 0] * (m[11] * m[13] - m[ 9] * m[15]) - + m[ 8] * (m[ 3] * m[13] - m[ 1] * m[15]) - + m[12] * (m[ 1] * m[11] - m[ 3] * m[ 9]); + + inv[10] = m[ 0] * (m[ 5] * m[15] - m[ 7] * m[13]) - + m[ 4] * (m[ 1] * m[15] - m[ 3] * m[13]) - + m[12] * (m[ 3] * m[ 5] - m[ 1] * m[ 7]); + + inv[11] = m[ 0] * (m[ 7] * m[ 9] - m[ 5] * m[11]) - + m[ 4] * (m[ 3] * m[ 9] - m[ 1] * m[11]) - + m[ 8] * (m[ 1] * m[ 7] - m[ 3] * m[ 5]); + + inv[12] = m[ 4] * (m[10] * m[13] - m[ 9] * m[14]) - + m[ 8] * (m[ 6] * m[13] - m[ 5] * m[14]) - + m[12] * (m[ 5] * m[10] - m[ 6] * m[ 9]); + + inv[13] = m[ 0] * (m[ 9] * m[14] - m[10] * m[13]) - + m[ 8] * (m[ 1] * m[14] - m[ 2] * m[13]) - + m[12] * (m[ 2] * m[ 9] - m[ 1] * m[10]); + + inv[14] = m[ 0] * (m[ 6] * m[13] - m[ 5] * m[14]) - + m[ 4] * (m[ 2] * m[13] - m[ 1] * m[14]) - + m[12] * (m[ 1] * m[ 6] - m[ 2] * m[ 5]); + + inv[15] = m[ 0] * (m[ 5] * m[10] - m[ 6] * m[ 9]) - + m[ 4] * (m[ 1] * m[10] - m[ 2] * m[ 9]) - + m[ 8] * (m[ 2] * m[ 5] - m[ 1] * m[ 6]); + + T det = m[0] * inv[ 0] + + m[1] * inv[ 4] + + m[2] * inv[ 8] + + m[3] * inv[12]; + + if (det == T (0.0)) return; + + det = T (1.0) / det; + + for (int i = 0; i < 16; ++i) + inv[i] *= det; +} + +// ======================================================================= +// function : GenerateCornerRays +// purpose : Generates primary rays for corners of screen quad +// ======================================================================= +void GenerateCornerRays (const GLdouble theInvModelProj[16], + float theOrigins[16], + float theDirects[16]) +{ + int aOriginIndex = 0; + int aDirectIndex = 0; + + for (int y = -1; y <= 1; y += 2) + { + for (int x = -1; x <= 1; x += 2) + { + OpenGl_RTVec4f aOrigin (float(x), + float(y), + -1.f, + 1.f); + + aOrigin = MatVecMult (theInvModelProj, aOrigin); + + OpenGl_RTVec4f aDirect (float(x), + float(y), + 1.f, + 1.f); + + aDirect = MatVecMult (theInvModelProj, aDirect) - aOrigin; + + GLdouble aInvLen = 1.f / sqrt (aDirect.x() * aDirect.x() + + aDirect.y() * aDirect.y() + + aDirect.z() * aDirect.z()); + + theOrigins [aOriginIndex++] = static_cast (aOrigin.x()); + theOrigins [aOriginIndex++] = static_cast (aOrigin.y()); + theOrigins [aOriginIndex++] = static_cast (aOrigin.z()); + theOrigins [aOriginIndex++] = 1.f; + + theDirects [aDirectIndex++] = static_cast (aDirect.x() * aInvLen); + theDirects [aDirectIndex++] = static_cast (aDirect.y() * aInvLen); + theDirects [aDirectIndex++] = static_cast (aDirect.z() * aInvLen); + theDirects [aDirectIndex++] = 0.f; + } + } +} + +// ======================================================================= +// function : Raytrace +// purpose : Redraws the window using OpenCL ray tracing +// ======================================================================= +Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView, + const int theSizeX, + int theSizeY, + const Tint theToSwap) +{ + if (!InitOpenCL()) + return Standard_False; + + if (!ResizeRaytraceOutputBuffer (theSizeX, theSizeY)) + return Standard_False; + + if (!UpdateRaytraceEnvironmentMap()) + return Standard_False; + + if (!UpdateRaytraceGeometry (Standard_True)) + return Standard_False; + + // Get model-view and projection matrices + TColStd_Array2OfReal theOrientation (0, 3, 0, 3); + TColStd_Array2OfReal theViewMapping (0, 3, 0, 3); + + myView->GetMatrices (theOrientation, theViewMapping, Standard_True); + + GLdouble aOrientationMatrix[16]; + GLdouble aViewMappingMatrix[16]; + GLdouble aOrientationInvers[16]; + + for (int j = 0; j < 4; ++j) + for (int i = 0; i < 4; ++i) + { + aOrientationMatrix [4 * j + i] = theOrientation (i, j); + aViewMappingMatrix [4 * j + i] = theViewMapping (i, j); + } + + ComputeInverseMatrix (aOrientationMatrix, aOrientationInvers); + + if (!UpdateRaytraceLightSources (aOrientationInvers)) + return Standard_False; + + // Generate primary rays for corners of the screen quad + glMatrixMode (GL_MODELVIEW); + + glLoadMatrixd (aViewMappingMatrix); + glMultMatrixd (aOrientationMatrix); + + GLdouble aModelProject[16]; + GLdouble aInvModelProj[16]; + + glGetDoublev (GL_MODELVIEW_MATRIX, aModelProject); + + ComputeInverseMatrix (aModelProject, aInvModelProj); + + GLfloat aOrigins[16]; + GLfloat aDirects[16]; + + GenerateCornerRays (aInvModelProj, + aOrigins, + aDirects); + + // Compute ray-traced image using OpenCL kernel + cl_mem anImages[] = { myRaytraceOutputImage, myRaytraceOutputImageSmooth }; + cl_int anError = clEnqueueAcquireGLObjects (myRaytraceQueue, + 2, anImages, + 0, NULL, NULL); + clFinish (myRaytraceQueue); + + if (myIsRaytraceDataValid) + { + RunRaytraceOpenCLKernels (theCView, + aOrigins, + aDirects, + theSizeX, + theSizeY); + } + + anError |= clEnqueueReleaseGLObjects (myRaytraceQueue, + 2, anImages, + 0, NULL, NULL); + clFinish (myRaytraceQueue); + + // Draw background + glPushAttrib (GL_ENABLE_BIT | + GL_CURRENT_BIT | + GL_COLOR_BUFFER_BIT | + GL_DEPTH_BUFFER_BIT); + + glDisable (GL_DEPTH_TEST); + + if (NamedStatus & OPENGL_NS_WHITEBACK) + { + glClearColor (1.0f, 1.0f, 1.0f, 1.0f); + } + else + { + glClearColor (myBgColor.rgb[0], + myBgColor.rgb[1], + myBgColor.rgb[2], + 1.0f); + } + + glClear (GL_COLOR_BUFFER_BIT); + + Handle(OpenGl_Workspace) aWorkspace (this); + myView->DrawBackground (aWorkspace); + + // Draw dummy quad to show result image + glEnable (GL_COLOR_MATERIAL); + glEnable (GL_BLEND); + + glDisable (GL_DEPTH_TEST); + + glBlendFunc (GL_ONE, GL_SRC_ALPHA); + + glEnable (GL_TEXTURE_RECTANGLE); + + glMatrixMode (GL_PROJECTION); + glLoadIdentity(); + + glMatrixMode (GL_MODELVIEW); + glLoadIdentity(); + + glColor3f (1.0f, 1.0f, 1.0f); + + glBindTexture (GL_TEXTURE_RECTANGLE, myRaytraceOutputTexture[theCView.IsAntialiasingEnabled ? 1 : 0]); + + if (myIsRaytraceDataValid) + { + glBegin (GL_QUADS); + { + glTexCoord2i ( 0, 0); glVertex2f (-1.f, -1.f); + glTexCoord2i ( 0, theSizeY); glVertex2f (-1.f, 1.f); + glTexCoord2i (theSizeX, theSizeY); glVertex2f ( 1.f, 1.f); + glTexCoord2i (theSizeX, 0); glVertex2f ( 1.f, -1.f); + } + glEnd(); + } + + glPopAttrib(); + + // Swap the buffers + if (theToSwap) + { + GetGlContext()->SwapBuffers(); + myBackBufferRestored = Standard_False; + } + else + glFlush(); + + return Standard_True; +} + +#endif diff --git a/src/TKOpenGl/EXTERNLIB b/src/TKOpenGl/EXTERNLIB index 0ead3dee1d..058f5720ad 100755 --- a/src/TKOpenGl/EXTERNLIB +++ b/src/TKOpenGl/EXTERNLIB @@ -5,6 +5,7 @@ CSF_OpenGlLibs CSF_objc CSF_Appkit CSF_IOKit +CSF_OPENCL CSF_FREETYPE CSF_GL2PS CSF_user32 diff --git a/src/V3d/V3d_View.cdl b/src/V3d/V3d_View.cdl index e9543ea47c..a21d80679a 100755 --- a/src/V3d/V3d_View.cdl +++ b/src/V3d/V3d_View.cdl @@ -1580,6 +1580,38 @@ is ---Purpose: Get clip planes. -- @return sequence clip planes that have been set for the view + SetRaytracingMode (me : mutable) is static; + ---Level: Public + ---Purpose: enables OpenCL-based ray-tracing mode + + SetRasterizationMode (me : mutable) is static; + ---Level: Public + ---Purpose: enables OpenGL-based rasterization mode + + EnableRaytracedShadows (me : mutable) is static; + ---Level: Public + ---Purpose: enables sharp shadows in OpenCL-based ray-tracing mode + + EnableRaytracedReflections (me : mutable) is static; + ---Level: Public + ---Purpose: enables specular reflections in OpenCL-based ray-tracing mode + + EnableRaytracedAntialiasing (me : mutable) is static; + ---Level: Public + ---Purpose: enables antialiasing in OpenCL-based ray-tracing mode + + DisableRaytracedShadows (me : mutable) is static; + ---Level: Public + ---Purpose: disables sharp shadows in OpenCL-based ray-tracing mode + + DisableRaytracedReflections (me : mutable) is static; + ---Level: Public + ---Purpose: disables specular reflections in OpenCL-based ray-tracing mode + + DisableRaytracedAntialiasing (me : mutable) is static; + ---Level: Public + ---Purpose: disables antialiasing in OpenCL-based ray-tracing mode + fields MyType : TypeOfView from V3d is protected ; diff --git a/src/V3d/V3d_View_5.cxx b/src/V3d/V3d_View_5.cxx index 06212e9a8a..5498d2e614 100755 --- a/src/V3d/V3d_View_5.cxx +++ b/src/V3d/V3d_View_5.cxx @@ -57,6 +57,62 @@ Standard_Boolean V3d_View::IsGLLightEnabled() const return MyView->IsGLLightEnabled(); } +void V3d_View::SetRaytracingMode() +{ + Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView(); + + cView->IsRaytracing = 1; +} + +void V3d_View::SetRasterizationMode() +{ + Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView(); + + cView->IsRaytracing = 0; +} + +void V3d_View::EnableRaytracedShadows() +{ + Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView(); + + cView->IsShadowsEnabled = 1; +} + +void V3d_View::EnableRaytracedReflections() +{ + Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView(); + + cView->IsReflectionsEnabled = 1; +} + +void V3d_View::EnableRaytracedAntialiasing() +{ + Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView(); + + cView->IsAntialiasingEnabled = 1; +} + +void V3d_View::DisableRaytracedShadows() +{ + Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView(); + + cView->IsShadowsEnabled = 0; +} + +void V3d_View::DisableRaytracedReflections() +{ + Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView(); + + cView->IsReflectionsEnabled = 0; +} + +void V3d_View::DisableRaytracedAntialiasing() +{ + Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView(); + + cView->IsAntialiasingEnabled = 0; +} + void V3d_View::SetLayerMgr(const Handle(V3d_LayerMgr)& aMgr) { MyLayerMgr = aMgr; diff --git a/src/ViewerTest/ViewerTest_ViewerCommands.cxx b/src/ViewerTest/ViewerTest_ViewerCommands.cxx index c8853270aa..e32241fe0e 100755 --- a/src/ViewerTest/ViewerTest_ViewerCommands.cxx +++ b/src/ViewerTest/ViewerTest_ViewerCommands.cxx @@ -19,7 +19,7 @@ // and conditions governing the rights and limitations under the License. #ifdef HAVE_CONFIG_H -# include + #include #endif #include @@ -5275,6 +5275,183 @@ static int VSetTextureMode (Draw_Interpretor& theDi, Standard_Integer theArgsNb, return 0; } +//============================================================================== +//function : VClInfo +//purpose : Prints info about active OpenCL device +//============================================================================== + +static Standard_Integer VClInfo (Draw_Interpretor& theInterpretor, + Standard_Integer, + const char**) +{ +#ifndef HAVE_OPENCL + + theInterpretor << "OCCT was compiled without OpenCL support!\n"; + +#else + + Handle(AIS_InteractiveContext) aContextAIS = ViewerTest::GetAISContext(); + + if (aContextAIS.IsNull()) + { + theInterpretor << "Call vinit before!\n"; + return 1; + } + + Handle(OpenGl_GraphicDriver) aDrv = Handle(OpenGl_GraphicDriver)::DownCast (aContextAIS->CurrentViewer()->Driver()); + + Graphic3d_CView* aCView = (Graphic3d_CView*) ViewerTest::CurrentView()->View()->CView(); + + NCollection_DataMap anInfo; + + if (aDrv.IsNull() || aCView == NULL || !aDrv->GetOpenClDeviceInfo (*aCView, anInfo)) + { + theInterpretor << "Cannot get OpenCL device info!\n"; + return 0; + } + + theInterpretor << "OpenCL device info:\n"; + + NCollection_DataMap::Iterator anIter (anInfo); + + for (; anIter.More(); anIter.Next()) + { + theInterpretor << anIter.Key() << ": \t" << anIter.Value() << "\n"; + } + +#endif + + return 0; +} + +//======================================================================= +//function : VRaytrace +//purpose : Enables/disables OpenCL-based ray-tracing +//======================================================================= + +#ifndef HAVE_OPENCL + +static Standard_Integer VRaytrace (Draw_Interpretor& theInterpretor, + Standard_Integer, + const char**) +{ + theInterpretor << "OCCT was compiled without OpenCL support!\n"; + + return 0; +} + +#else + +static Standard_Integer VRaytrace (Draw_Interpretor&, + Standard_Integer theArgNb, + const char** theArgVec) +{ + Handle(AIS_InteractiveContext) aContext = ViewerTest::GetAISContext(); + + if (aContext.IsNull()) + { + std::cerr << "Use 'vinit' command before " << theArgVec[0] << "\n"; + return 1; + } + + if (theArgNb < 2) + { + std::cerr << "Usage : " << theArgVec[0] << " 0|1\n"; + return 1; + } + + Standard_Integer isOn = atoi(theArgVec[1]); + + Handle(V3d_View) aView = ViewerTest::CurrentView(); + + if (isOn) + aView->SetRaytracingMode(); + else + aView->SetRasterizationMode(); + + aView->Redraw(); + + return 0; +} + +#endif + +//======================================================================= +//function : VSetRaytraceMode +//purpose : Enables/disables features of OpenCL-based ray-tracing +//======================================================================= + +#ifndef HAVE_OPENCL + +static Standard_Integer VSetRaytraceMode (Draw_Interpretor& theInterpretor, + Standard_Integer, + const char**) +{ + theInterpretor << "OCCT was compiled without OpenCL support!\n"; + + return 0; +} + +#else + +static Standard_Integer VSetRaytraceMode (Draw_Interpretor&, + Standard_Integer theArgNb, + const char ** theArgVec) +{ + Handle(AIS_InteractiveContext) aContext = ViewerTest::GetAISContext(); + + if (aContext.IsNull()) + { + std::cerr << "Use 'vinit' command before " << theArgVec[0] << "\n"; + return 1; + } + + if (theArgNb < 2) + { + std::cerr << "Usage : " << theArgVec[0] << " [shad=0|1] [refl=0|1] [aa=0|1]\n"; + return 1; + } + + Handle(V3d_View) aView = ViewerTest::CurrentView(); + + for (Standard_Integer anArgIter = 1; anArgIter < theArgNb; ++anArgIter) + { + const TCollection_AsciiString anArg (theArgVec[anArgIter]); + + if (anArg.Search ("shad=") > -1) + { + if (anArg.Token ("=", 2).IntegerValue() != 0) + aView->EnableRaytracedShadows(); + else + aView->DisableRaytracedShadows(); + } + else if (anArg.Search ("refl=") > -1) + { + if (anArg.Token ("=", 2).IntegerValue() != 0) + aView->EnableRaytracedReflections(); + else + aView->DisableRaytracedReflections(); + } + else if (anArg.Search ("aa=") > -1) + { + if (anArg.Token ("=", 2).IntegerValue() != 0) + aView->EnableRaytracedAntialiasing(); + else + aView->DisableRaytracedAntialiasing(); + } + else + { + std::cerr << "Unknown argument: " << anArg << "\n"; + } + } + + aView->Redraw(); + + return 0; +} + +#endif + //======================================================================= //function : ViewerCommands //purpose : @@ -5548,4 +5725,13 @@ void ViewerTest::ViewerCommands(Draw_Interpretor& theCommands) " 2 - all textures enabled.\n" " this command sets texture details mode for the specified view.\n" , __FILE__, VSetTextureMode, group); + theCommands.Add("vraytrace", + "vraytrace 0|1", + __FILE__,VRaytrace,group); + theCommands.Add("vclinfo", + "vclinfo", + __FILE__,VClInfo,group); + theCommands.Add("vsetraytracemode", + "vsetraytracemode [shad=0|1] [refl=0|1] [aa=0|1]", + __FILE__,VSetRaytraceMode,group); } diff --git a/tests/bugs/parse.rules b/tests/bugs/parse.rules index eab3eb5026..466a2c0b25 100755 --- a/tests/bugs/parse.rules +++ b/tests/bugs/parse.rules @@ -1,4 +1,5 @@ FAILED /\bFaulty\b/ bad shape IGNORE /^Error [23]d = [\d.-]+/ debug output of blend command SKIPPED /Error: unsupported locale specification/ locale is unavailable on tested system +SKIPPED /OCCT was compiled without OpenCL support!/ OK /Relative error of mass computation/ message from vprops diff --git a/tests/bugs/vis/bug24130 b/tests/bugs/vis/bug24130 new file mode 100644 index 0000000000..7c30ba3834 --- /dev/null +++ b/tests/bugs/vis/bug24130 @@ -0,0 +1,48 @@ +puts "TODO OCC24130 Debian60-64: OCCT was compiled without OpenCL support!" +puts "TODO OCC24130 Windows: TKOpenGl | Type\: Error | ID\: 0 | Severity\: High | Message\:" + +puts "========" +puts "OCC24130 Implementing ray tracing visualization core" +puts "========" + +# custom shapes +set aShape1 [locate_data_file occ/Top.brep] +set aShape2 [locate_data_file occ/Bottom.brep] + +# setup 3D viewer content +vinit name=View1 w=512 h=512 +vglinfo + +vvbo 0 +vsetdispmode 1 +vsetgradientbg 180 200 255 180 180 180 2 +restore $aShape1 s1 +restore $aShape2 s2 +vdisplay s1 s2 +vsetmaterial s1 Silver +vsetmaterial s2 Pewter +vfit + +# activate ray-tracing +vraytrace 1 +vclinfo + +set aModeNum 0 +for { set aAAMode 0 } { $aAAMode <= 1 } { incr aAAMode } { + for { set aReflMode 0 } { $aReflMode <= 1 } { incr aReflMode } { + for { set aShadMode 0 } { $aShadMode <= 1 } { incr aShadMode } { + vsetraytracemode shad=$aShadMode refl=$aReflMode aa=$aAAMode + vdump $imagedir/${casename}_${aModeNum}.png + incr aModeNum + } + } +} + +vtextureenv on 5 +for { set aAAMode 0 } { $aAAMode <= 1 } { incr aAAMode } { + for { set aShadMode 0 } { $aShadMode <= 1 } { incr aShadMode } { + vsetraytracemode shad=$aShadMode refl=1 aa=$aAAMode + vdump $imagedir/${casename}_${aModeNum}.png + incr aModeNum + } +} -- 2.20.1