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.
*.jxx eol=lf
*.lxx eol=lf
*.pxx eol=lf
+*.cl eol=lf
*.cdl eol=lf
*.edl eol=lf
*.yacc eol=lf
p BOPDS
p BOPCol
p BOPInt
+r OpenCL
#include <QApplication>
#include <QSignalMapper>
+#include <Graphic3d_GraphicDriver.hxx>
+#include <OpenGl_GraphicDriver.hxx>
+
#include <stdlib.h>
static ApplicationCommonWindow* stApp;
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") );
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") );
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
view->addAction( viewStatusAction );
// add a help menu
-
QMenu * help = new QMenu( this );
menuBar()->addSeparator();
help = menuBar()->addMenu( QObject::tr("MNU_HELP") );
myStdToolBar->addAction( helpAboutAction );
myStdActions.at(FileCloseId)->setEnabled(myDocuments.count() > 0);
+
+ myStdActions.at(FilePrefUseVBOId)->setEnabled( true );
}
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 ) ) );
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" ) );
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 )
{
}
}
+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 );
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()
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();
#include <QWorkspace>
#include <QList>
+
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();
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*/ );
QList<QAction*> myStdActions;
QList<QAction*> myToolActions;
+ QList<QAction*> myRaytraceActions;
QList<QAction*> myMaterialActions;
QList<DocumentCommon*> myDocuments;
QToolBar* myStdToolBar;
QToolBar* myCasCadeBar;
+ QToolBar* myRaytraceBar;
QMenu* myFilePopup;
QMenu* myWindowPopup;
QAction* myFileSeparator;
<source>ICON_NEW</source>
<translation>new.png</translation>
</message>
+ <message>
+ <source>ICON_NEW_GL</source>
+ <translation>newGL.png</translation>
+ </message>
+ <message>
+ <source>ICON_NEW_RT</source>
+ <translation>newRT.png</translation>
+ </message>
<message>
<source>ICON_VIEW_RIGHT</source>
<translation>view_right.png</translation>
<source>ICON_SAMPLE</source>
<translation>lamp.png</translation>
</message>
+ <message>
+ <source>ICON_TOOL_SHADOWS</source>
+ <translation>shadows.png</translation>
+ </message>
+ <message>
+ <source>ICON_TOOL_REFLECTIONS</source>
+ <translation>reflections.png</translation>
+ </message>
+ <message>
+ <source>ICON_TOOL_ANTIALIASING</source>
+ <translation>antialiasing.png</translation>
+ </message>
</context>
</TS>
<source>MNU_FILE</source>
<translation>&File</translation>
</message>
+ <message>
+ <source>MNU_PREFERENCES</source>
+ <translation>&Preferences</translation>
+ </message>
+ <message>
+ <source>MNU_USE_VBO</source>
+ <translation>&Use VBO</translation>
+ </message>
<message>
<source>MNU_GOLD</source>
<translation>&Gold</translation>
<source>MNU_TOOL_TRANS</source>
<translation>&Transpatency</translation>
</message>
+ <message>
+ <source>MNU_TOOL_SHADOWS</source>
+ <translation>&Shadows</translation>
+ </message>
+ <message>
+ <source>MNU_TOOL_REFLECTIONS</source>
+ <translation>&Reflections</translation>
+ </message>
+ <message>
+ <source>MNU_TOOL_ANTIALIASING</source>
+ <translation>&Antialiasing</translation>
+ </message>
<message>
<source>BTN_BRASS</source>
<translation>Brass</translation>
<source>MNU_CH_BACK</source>
<translation>&Change Background</translation>
</message>
+ <message>
+ <source>MNU_CH_ENV_MAP</source>
+ <translation>&Environment Map</translation>
+ </message>
<message>
<source>TBR_CH_BACK</source>
<translation>Change Background</translation>
</message>
<message>
<source>TBR_WINDOW_NEW3D</source>
- <translation>New 3d View</translation>
+ <translation>New 3D View</translation>
+ </message>
+ <message>
+ <source>TBR_WINDOW_NEW3D_GL</source>
+ <translation>New GL 3D View</translation>
+ </message>
+ <message>
+ <source>MNU_WINDOW_NEW3D_GL</source>
+ <translation>New GL 3D View</translation>
+ </message>
+ <message>
+ <source>TBR_WINDOW_NEW3D_RT</source>
+ <translation>&New RT 3D View</translation>
</message>
<message>
- <source>MNU_WINDOW_NEW3D</source>
- <translation>&New 3d View</translation>
+ <source>MNU_WINDOW_NEW3D_RT</source>
+ <translation>&New RT 3D View</translation>
</message>
<message>
<source>MNU_STATUS_BAR</source>
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()
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;
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
void fitAll();
protected:
- virtual MDIWindow* createNewMDIWindow();
+ virtual MDIWindow* createNewMDIWindow( bool theRT = false );
signals:
void selectionChanged();
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();
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 );
setCentralWidget( vb );
myDocument = aDocument;
- myView = new View( myDocument->getContext(), vb );
+ myView = new View( myDocument->getContext(), vb, theRT );
layout->addWidget( myView );
connect( myView, SIGNAL( selectionChanged() ),
resize( sizeHint() );
setFocusPolicy( Qt::StrongFocus );
+
+#ifdef HAVE_OPENCL
+
+ myShadowsEnabled = true;
+ myReflectionsEnabled = true;
+ myAntialiasingEnabled = false;
+
+#endif
}
MDIWindow::~MDIWindow()
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
+
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();
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 );
protected:
DocumentCommon* myDocument;
View* myView;
+
+#ifdef HAVE_OPENCL
+
+ bool myShadowsEnabled;
+ bool myReflectionsEnabled;
+ bool myAntialiasingEnabled;
+
+#endif
};
#endif
#include <QColorDialog>
#include <QCursor>
#include <QFileInfo>
+#include <QFileDialog>
#include <QMouseEvent>
#include <QRubberBand>
#include <Visual3d_View.hxx>
#include <Graphic3d_ExportFormat.hxx>
#include <Graphic3d_GraphicDriver.hxx>
+#include <Graphic3d_TextureEnv.hxx>
#include <QWindowsStyle>
#if defined(_WIN32) || defined(__WIN32__)
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();
myFirst = true;
myContext = theContext;
+ //if (theRT)
+ // myContext->SetDisplayMode( AIS_DisplayMode::AIS_Shaded, 1 );
+
myXmin = 0;
myYmin = 0;
myXmax = 0;
View::~View()
{
+ delete myBackMenu;
}
void View::init()
}
myView->SetBackgroundColor (Quantity_NOC_BLACK);
myView->MustBeResized();
+
+ if (myIsRT)
+ myView->SetRaytracingMode();
}
void View::paintEvent( QPaintEvent * )
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();
}
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();
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();
ViewAxoId, ViewRotationId, ViewResetId, ViewHlrOffId, ViewHlrOnId };
View( Handle(AIS_InteractiveContext) theContext,
- QWidget* parent);
+ QWidget* parent,
+ bool theRT = false );
+
~View();
virtual void init();
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 );
void hlrOff();
void updateToggled( bool );
void onBackground();
+ void onEnvironmentMap();
protected:
virtual void paintEvent( QPaintEvent* );
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;
Quantity_Factor myCurZoom;
Standard_Boolean myHlrModeIsOn;
QList<QAction*>* myViewActions;
+ QMenu* myBackMenu;
QRubberBand* myRectBand; //!< selection rectangle rubber band
};
GlobalSection(SolutionProperties) = preSolution
HideSolutionNode = FALSE
EndGlobalSection
-EndGlobal
+EndGlobal
\ No newline at end of file
Backfacing (0),
GDisplayCB (NULL),
GClientData (NULL),
- ptrFBO (NULL)
+ ptrFBO (NULL),
+ WasRedrawnGL (0),
+ IsRaytracing (0),
+ IsShadowsEnabled (1),
+ IsReflectionsEnabled (1),
+ IsAntialiasingEnabled (0)
{
//
}
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);
CSF_Appkit
CSF_IOKit
CSF_OpenGlLibs
+CSF_OPENCL
CSF_AviLibs
CSF_FREETYPE
CSF_GL2PS
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
--- /dev/null
+// 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 <config.h>
+#endif
+
+#ifdef HAVE_OPENCL
+
+#include <Standard_ShortReal.hxx>
+
+#include <OpenGl_AABB.hxx>
+
+
+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
--- /dev/null
+// 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 <OpenGl_RaytraceTypes.hxx>
+
+
+//! 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
#else
contextDebug (Standard_False),
#endif
- contextNoAccel (Standard_False)
+ contextNoAccel (Standard_False),
+ keepArrayData (Standard_False)
{
//
}
*/
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.
--- /dev/null
+// 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 <OpenGl_GlCore20.hxx>
+
+#if defined(__APPLE__) || defined(__MACOSX)
+ #include <OpenCL/opencl.h>
+#else
+ #include <CL/cl.h>
+ #include <CL/cl_gl.h>
+#endif
+
+#endif // _OpenGl_Cl_H__
#include <OpenGl_Display.hxx>
#ifdef HAVE_CONFIG_H
-# include <config.h>
+ #include <config.h>
#endif
#ifdef HAVE_GL2PS
// 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 <config.h>
+#endif
+
#include <OpenGl_GraphicDriver.hxx>
#include <OpenGl_Context.hxx>
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<TCollection_AsciiString, TCollection_AsciiString>&)
+{
+ return Standard_True;
+}
+
+#else
+
+Standard_Boolean OpenGl_GraphicDriver::GetOpenClDeviceInfo (const Graphic3d_CView& theCView,
+ NCollection_DataMap<TCollection_AsciiString, TCollection_AsciiString>& theInfo)
+{
+
+ if (theCView.ViewId == -1 || theCView.ptrView == NULL)
+ {
+ return Standard_False;
+ }
+
+ return reinterpret_cast<const OpenGl_CView*> (theCView.ptrView)->WS->GetOpenClDeviceInfo (theInfo);
+}
+
+#endif
+
// =======================================================================
// function : BeginAddMode
// purpose :
//! 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<TCollection_AsciiString, TCollection_AsciiString>& theInfo);
+
private:
//! Method to retrieve valid GL context.
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)
{
// purpose or non-infringement. Please see the License for the specific terms
// and conditions governing the rights and limitations under the License.
-#include <OpenGl_Group.hxx>
+#ifdef HAVE_CONFIG_H
+ #include <config.h>
+#endif
+
+#include <OpenGl_Group.hxx>
#include <OpenGl_PrimitiveArray.hxx>
+#include <OpenGl_Structure.hxx>
#include <OpenGl_Workspace.hxx>
// =======================================================================
// 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
}
// =======================================================================
anAspectFace->SetAspect (theAspect);
AddElement (TelNil/*TelAspectFace*/, anAspectFace);
}
+
+#ifdef HAVE_OPENCL
+ if (myIsRaytracable)
+ {
+ myModificationState++;
+
+ if (myAncestorStructure != NULL)
+ {
+ myAncestorStructure->UpdateStateWithAncestorStructures();
+ }
+ }
+#endif
}
// =======================================================================
// 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
}
// =======================================================================
#include <OpenGl_tsm.hxx>
class OpenGl_Group;
+class OpenGl_Structure;
typedef NCollection_List<const OpenGl_Group* > OpenGl_ListOfGroup;
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);
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:
// and conditions governing the rights and limitations under the License.
+#ifdef HAVE_CONFIG_H
+ #include <config.h>
+#endif
+
#include <OpenGl_GlCore11.hxx>
#include <OpenGl_LayerList.hxx>
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
}
//=======================================================================
if (aList.Remove (theStructure) >= 0)
{
myNbStructures--;
+
+#ifdef HAVE_OPENCL
+ if (theStructure->IsRaytracable())
+ {
+ myModificationState++;
+ }
+#endif
+
return;
}
if (aScanList.Remove (theStructure) >= 0)
{
myNbStructures--;
+
+#ifdef HAVE_OPENCL
+ if (theStructure->IsRaytracable())
+ {
+ myModificationState++;
+ }
+#endif
+
return;
}
}
//! 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
Standard_Integer myNbPriorities;
Standard_Integer myNbStructures;
+#ifdef HAVE_OPENCL
+ mutable Standard_Size myModificationState;
+#endif
+
public:
DEFINE_STANDARD_ALLOC
};
}
}
- clearMemoryOwn();
+ if (!aGlCtx->caps->keepArrayData)
+ {
+ clearMemoryOwn();
+ }
+
return Standard_True;
}
//! 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;
--- /dev/null
+// 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 <config.h>
+#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
--- /dev/null
+// 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 <vector>
+
+#include <NCollection_Vec4.hxx>
+#include <NCollection_StdAllocator.hxx>
+
+//! 4D vector of integers.
+typedef NCollection_Vec4<int> OpenGl_RTVec4i;
+
+//! 4D vector of floats.
+typedef NCollection_Vec4<float> OpenGl_RTVec4f;
+
+//! 4D vector of doubles.
+typedef NCollection_Vec4<double> OpenGl_RTVec4d;
+
+//! Array of 4D integer vectors.
+typedef std::vector<OpenGl_RTVec4i,
+ NCollection_StdAllocator<OpenGl_RTVec4i> > OpenGl_RTArray4i;
+
+//! Array of 4D floating point vectors.
+typedef std::vector<OpenGl_RTVec4f,
+ NCollection_StdAllocator<OpenGl_RTVec4f> > OpenGl_RTArray4f;
+
+#endif
--- /dev/null
+// 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 <config.h>
+#endif
+
+#ifdef HAVE_OPENCL
+
+#include <limits>
+
+#include <OpenGl_SceneGeometry.hxx>
+
+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<OpenGl_RaytraceMaterial,
+ NCollection_StdAllocator<OpenGl_RaytraceMaterial> > 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<int> (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<int> (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 <iostream>
+#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<int> (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<int> (theGeometry.Triangles.size() - 1)));
+
+ // Building nodes while tasks queue is not empty
+ for (int aTaskId = 0; aTaskId < static_cast<int> (myNodeTasksQueue.size()); ++aTaskId)
+ {
+ BuildNode (theGeometry, aTaskId);
+ }
+
+ // Write support data to optimize traverse
+ for (int aNode = 0; aNode < static_cast<int> (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<float>::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
--- /dev/null
+// 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 <OpenGl_AABB.hxx>
+#include <OpenGl_Structure.hxx>
+#include <OpenGl_PrimitiveArray.hxx>
+
+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<float*> (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<float*> (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<OpenGl_RaytraceMaterial,
+ NCollection_StdAllocator<OpenGl_RaytraceMaterial> > Materials;
+
+ //! Array of properties of light sources.
+ std::vector<OpenGl_RaytraceLight,
+ NCollection_StdAllocator<OpenGl_RaytraceLight> > 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_BVHBin,
+ NCollection_StdAllocator<OpenGl_BVHBin> > 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<OpenGl_BVHNodeTask> myNodeTasksQueue;
+
+ //! Builded BVH tree.
+ OpenGl_BVH myTree;
+
+ //! Maximum depth of BVH tree.
+ int myMaxDepth;
+};
+
+#endif
+#endif
// 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 <config.h>
+#endif
+
#include <OpenGl_CappingAlgo.hxx>
#include <OpenGl_Context.hxx>
myNamedStatus(0),
myZLayer(0)
{
+#if HAVE_OPENCL
+ myIsRaytracable = Standard_False;
+ myModificationState = 0;
+#endif
}
// =======================================================================
// 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
}
// =======================================================================
myAspectFace = new OpenGl_AspectFace();
}
myAspectFace->SetAspect (theAspect);
+
+#ifdef HAVE_OPENCL
+ if (myIsRaytracable)
+ {
+ UpdateStateWithAncestorStructures();
+ }
+#endif
}
// =======================================================================
}
else
{
+#ifndef HAVE_OPENCL
myHighlightBox = new OpenGl_Group();
+#else
+ myHighlightBox = new OpenGl_Group (this);
+#endif
}
CALL_DEF_CONTEXTLINE aContextLine;
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();
// 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;
}
// Check for the given group
if (anIter.Value() == theGroup)
{
- // Delete object
- OpenGl_Element::Destroy (theGlCtx, const_cast<OpenGl_Group*& > (anIter.ChangeValue()));
myGroups.Remove (anIter);
+
+#ifdef HAVE_OPENCL
+ if (theGroup->IsRaytracable())
+ {
+ UpdateStateWithAncestorStructures();
+ UpdateRaytracableWithAncestorStructures();
+ }
+#endif
+
+ // Delete object
+ OpenGl_Element::Destroy (theGlCtx, const_cast<OpenGl_Group*& > (theGroup));
return;
}
}
// =======================================================================
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<OpenGl_Group*& > (anIter.ChangeValue()));
}
myGroups.Clear();
+
+#ifdef HAVE_OPENCL
+ if (aRaytracableGroupDeleted)
+ {
+ UpdateStateWithAncestorStructures();
+ UpdateRaytracableWithAncestorStructures();
+ }
+#endif
}
// =======================================================================
#include <OpenGl_Group.hxx>
#include <OpenGl_Matrix.hxx>
+#include <OpenGl_NamedStatus.hxx>
#include <Graphic3d_SetOfHClipPlane.hxx>
class OpenGl_Structure : public OpenGl_Element
{
+ friend class OpenGl_Group;
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; }
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
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
// 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 <config.h>
+#endif
+
#include <NCollection_Mat4.hxx>
#include <OpenGl_Context.hxx>
myIntShadingMethod = TEL_SM_FLAT;
break;
}
+
+#ifdef HAVE_OPENCL
+ myModificationState = 1; // initial state
+#endif
}
/*----------------------------------------------------------------------*/
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
}
/*----------------------------------------------------------------------*/
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; }
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
Standard_Boolean myViewMappingChanged;
Standard_Boolean myLightSourcesChanged;
+#ifdef HAVE_OPENCL
+ Standard_Size myModificationState;
+#endif
+
public:
DEFINE_STANDARD_ALLOC
};
/*----------------------------------------------------------------------*/
-//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();
// - 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*/
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
GLfloat texX_range = 1.F; // texture <s> coordinate
GLfloat texY_range = 1.F; // texture <t> coordinate
- // Set up for stretching or tiling
+ // Set up for stretching or tiling
GLfloat x_offset, y_offset;
if ( myBgTexture.Style == Aspect_FM_CENTERED )
{
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);
// 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;
// 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 <config.h>
+#endif
+
#include <OpenGl_GlCore15.hxx>
#include <InterfaceGraphic.hxx>
#include <OpenGl_Context.hxx>
#include <OpenGl_FrameBuffer.hxx>
#include <OpenGl_Texture.hxx>
+#include <OpenGl_View.hxx>
#include <OpenGl_Workspace.hxx>
#include <OpenGl_Element.hxx>
// 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
}
// =======================================================================
// =======================================================================
OpenGl_Workspace::~OpenGl_Workspace()
{
+#ifdef HAVE_OPENCL
+ ReleaseOpenCL();
+#endif
}
// =======================================================================
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)
{
#ifndef _OpenGl_Workspace_Header
#define _OpenGl_Workspace_Header
+#ifdef HAVE_OPENCL
+ #include <map>
+ #include <set>
+
+ #include <OpenGl_Cl.hxx>
+#endif
+
#include <Handle_OpenGl_Workspace.hxx>
#include <OpenGl_Window.hxx>
#include <OpenGl_Matrix.hxx>
#include <OpenGl_NamedStatus.hxx>
#include <OpenGl_PrinterContext.hxx>
+#ifdef HAVE_OPENCL
+ #include <OpenGl_SceneGeometry.hxx>
+#endif
#include <OpenGl_TextParam.hxx>
#include <OpenGl_RenderFilter.hxx>
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:
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<TCollection_AsciiString, TCollection_AsciiString>& 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<const OpenGl_Structure*>& 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<const OpenGl_Structure*, Standard_Size> myStructureStates;
+
+#endif // HAVE_OPENCL
+
protected: //! @name protected fields
Handle(OpenGl_PrinterContext) myPrintContext;
--- /dev/null
+// 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 <config.h>
+#endif
+
+#ifdef HAVE_OPENCL
+
+#if defined(_WIN32)
+
+ #include <windows.h>
+ #include <wingdi.h>
+
+ #pragma comment (lib, "DelayImp.lib")
+ #pragma comment (lib, "OpenCL.lib")
+
+#elif defined(__APPLE__) && !defined(MACOSX_USE_GLX)
+ #include <OpenGL/CGLCurrent.h>
+#else
+ #include <GL/glx.h>
+#endif
+
+#include <OpenGl_Context.hxx>
+#include <OpenGl_Texture.hxx>
+#include <OpenGl_View.hxx>
+#include <OpenGl_Workspace.hxx>
+
+using namespace OpenGl_Raytrace;
+
+//! Use this macro to output ray-tracing debug info
+// #define RAY_TRACE_PRINT_INFO
+
+#ifdef RAY_TRACE_PRINT_INFO
+ #include <OSD_Timer.hxx>
+#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<float> (m[ 0] * v.x() + m[ 4] * v.y() +
+ m[ 8] * v.z() + m[12] * v.w()),
+ static_cast<float> (m[ 1] * v.x() + m[ 5] * v.y() +
+ m[ 9] * v.z() + m[13] * v.w()),
+ static_cast<float> (m[ 2] * v.x() + m[ 6] * v.y() +
+ m[10] * v.z() + m[14] * v.w()),
+ static_cast<float> (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<const OpenGl_Structure*> 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<const OpenGl_Structure*, Standard_Size>::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<const OpenGl_Structure*, Standard_Size>::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<const OpenGl_Structure*>& 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<int> (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<int> (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<int> (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<OpenGl_AspectFace*> (aNode->elem);
+
+ if (anAspect != NULL)
+ {
+ aMatID = static_cast<int> (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<OpenGl_PrimitiveArray*> (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<int> (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<TCollection_AsciiString,
+ TCollection_AsciiString>& 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<float> (
+ myRaytraceSceneData.Materials.size()) * sizeof (OpenGl_RaytraceMaterial);
+
+ aMemUsed += static_cast<float> (
+ myRaytraceSceneData.Triangles.size() * sizeof (OpenGl_RTVec4i) +
+ myRaytraceSceneData.Vertices.size() * sizeof (OpenGl_RTVec4f) +
+ myRaytraceSceneData.Normals.size() * sizeof (OpenGl_RTVec4f));
+
+ aMemUsed += static_cast<float> (
+ 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<cl_int> (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 <typename T>
+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<GLfloat> (aOrigin.x());
+ theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.y());
+ theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.z());
+ theOrigins [aOriginIndex++] = 1.f;
+
+ theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.x() * aInvLen);
+ theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.y() * aInvLen);
+ theDirects [aDirectIndex++] = static_cast<GLfloat> (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
CSF_objc
CSF_Appkit
CSF_IOKit
+CSF_OPENCL
CSF_FREETYPE
CSF_GL2PS
CSF_user32
---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 ;
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;
// and conditions governing the rights and limitations under the License.
#ifdef HAVE_CONFIG_H
-# include <config.h>
+ #include <config.h>
#endif
#include <OpenGl_GlCore20.hxx>
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<TCollection_AsciiString, TCollection_AsciiString> 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<TCollection_AsciiString, TCollection_AsciiString>::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 :
" 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);
}
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
--- /dev/null
+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
+ }
+}
projects / occt.git / commitdiff
