0024130: Implementing ray tracing visualization core

The purpose of this functionality is to bring a basic ray-tracing solution to existing OCCT visualization toolkit (TKOpenGL).
Currently ray-tracing visualization core supports sharp shadows, specular reflections, transparency and adaptive anti-aliasing.
However, the basis for all ray-tracing algorithms is versatile, allowing you to add new ray-tracing features easily (such as ambient occlusion).
All ray-tracing computations are performed on the GPU using OpenCL framework, allowing real-time rendering performance.

It is important to note, that real-time ray-tracing is possible using high-performance GPUs with support of OpenCL 1.1 and higher (such as NVIDIA GeForce 660 or ATI/AMD Radeon 7850).
When using low-end GPUs (such as NVIDIA GeForce 640) the ray-tracing performance may slow down significantly.
Therefore, even with NVIDIA GeForce 640 you can render scenes with the millions of triangles. The support of OpenCL-enabled CPUs and integrated graphics cards is not guaranteed.

.gitattributes

@@ -12,6 +12,7 @@
 *.jxx       eol=lf
 *.lxx       eol=lf
 *.pxx       eol=lf
+*.cl        eol=lf
 *.cdl       eol=lf
 *.edl       eol=lf
 *.yacc      eol=lf

adm/UDLIST

@@ -488,3 +488,4 @@ p BOPAlgo
 p BOPDS
 p BOPCol
 p BOPInt
+r OpenCL

samples/qt/Common/src/ApplicationCommon.cxx

@@ -12,6 +12,9 @@
 #include <QApplication>
 #include <QSignalMapper>
 
+#include <Graphic3d_GraphicDriver.hxx>
+#include <OpenGl_GraphicDriver.hxx>
+
 #include <stdlib.h>
 
 static ApplicationCommonWindow* stApp;
@@ -60,9 +63,8 @@ void ApplicationCommonWindow::createStandardOperations()
   helpIcon = QPixmap( dir + QObject::tr("ICON_HELP") );
   closeIcon = QPixmap( dir + QObject::tr("ICON_CLOSE") );
 
-  QAction * fileNewAction, * fileCloseAction, * fileQuitAction,
+  QAction * fileNewAction, * fileCloseAction, * filePrefUseVBOAction,
-          * viewToolAction, * viewStatusAction,
+          * fileQuitAction, * viewToolAction, * viewStatusAction, * helpAboutAction;
-		  * helpAboutAction;
 
   fileNewAction = new QAction( newIcon, QObject::tr("MNU_NEW"), this );
   fileNewAction->setToolTip( QObject::tr("TBR_NEW") );
@@ -78,6 +80,14 @@ void ApplicationCommonWindow::createStandardOperations()
   connect( fileCloseAction, SIGNAL( triggered() ) , this, SLOT( onCloseWindow() ) );
   myStdActions.insert( FileCloseId, fileCloseAction );
 
+  filePrefUseVBOAction = new QAction( QObject::tr("MNU_USE_VBO"), this );
+  filePrefUseVBOAction->setToolTip( QObject::tr("TBR_USE_VBO") );
+  filePrefUseVBOAction->setStatusTip( QObject::tr("TBR_USE_VBO") );
+  filePrefUseVBOAction->setCheckable( true );
+  filePrefUseVBOAction->setChecked( true );
+  connect( filePrefUseVBOAction, SIGNAL( activated() ) , this, SLOT( onUseVBO() ) );
+  myStdActions.insert( FilePrefUseVBOId, filePrefUseVBOAction );
+
   fileQuitAction = new QAction( QObject::tr("MNU_QUIT"), this );
   fileQuitAction->setToolTip( QObject::tr("TBR_QUIT") );
   fileQuitAction->setStatusTip( QObject::tr("TBR_QUIT") );
@@ -108,12 +118,18 @@ void ApplicationCommonWindow::createStandardOperations()
   connect( helpAboutAction, SIGNAL( triggered() ) , this, SLOT( onAbout() ) );
   myStdActions.insert( HelpAboutId, helpAboutAction );
 
+  // create preferences menu
+  QMenu* aPrefMenu = new QMenu( QObject::tr("MNU_PREFERENCES") );
+  aPrefMenu->addAction( filePrefUseVBOAction );
+  
   // popuplate a menu with all actions
   myFilePopup = new QMenu( this );
   myFilePopup = menuBar()->addMenu( QObject::tr("MNU_FILE") );
   myFilePopup->addAction( fileNewAction );
   myFilePopup->addAction( fileCloseAction );
   myFileSeparator = myFilePopup->addSeparator();
+  myFilePopup->addMenu( aPrefMenu );
+  myFileSeparator = myFilePopup->addSeparator();
   myFilePopup->addAction( fileQuitAction );
     
 	// add a view menu
@@ -124,7 +140,6 @@ void ApplicationCommonWindow::createStandardOperations()
   view->addAction( viewStatusAction );
 	
 	// add a help menu
-
   QMenu * help = new QMenu( this );
   menuBar()->addSeparator();
   help = menuBar()->addMenu( QObject::tr("MNU_HELP") );
@@ -136,6 +151,8 @@ void ApplicationCommonWindow::createStandardOperations()
   myStdToolBar->addAction( helpAboutAction );
 	
   myStdActions.at(FileCloseId)->setEnabled(myDocuments.count() > 0);
+
+  myStdActions.at(FilePrefUseVBOId)->setEnabled( true );
 }
 
 void ApplicationCommonWindow::createCasCadeOperations()
@@ -181,9 +198,43 @@ void ApplicationCommonWindow::createCasCadeOperations()
   a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_DEL") ), QObject::tr("MNU_TOOL_DEL"), this );
   a->setToolTip( QObject::tr("TBR_TOOL_DEL") );
   a->setStatusTip( QObject::tr("TBR_TOOL_DEL") );
-  connect( a, SIGNAL( triggered() ) , this, SLOT( onToolAction() ) );
+  connect( a, SIGNAL( activated() ) , this, SLOT( onToolAction() ) );
   myToolActions.insert( ToolDeleteId, a );
 
+#ifdef HAVE_OPENCL
+
+  // populate a tool bar with some actions
+  myRaytraceBar = addToolBar( tr( "Ray-trace options" ) );
+
+  a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_SHADOWS") ), QObject::tr("MNU_TOOL_SHADOWS"), this );
+  a->setToolTip( QObject::tr("TBR_TOOL_SHADOWS") );
+  a->setStatusTip( QObject::tr("TBR_TOOL_SHADOWS") );
+  a->setCheckable( true );
+  a->setChecked( true );
+  connect( a, SIGNAL( activated() ) , this, SLOT( onRaytraceAction() ) );
+  myRaytraceActions.insert( ToolShadowsId, a );
+  myRaytraceBar->addAction( a );
+
+  a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_REFLECTIONS") ), QObject::tr("MNU_TOOL_REFLECTIONS"), this );
+  a->setToolTip( QObject::tr("TBR_TOOL_REFLECTIONS") );
+  a->setStatusTip( QObject::tr("TBR_TOOL_REFLECTIONS") );
+  a->setCheckable( true );
+  a->setChecked( true );
+  connect( a, SIGNAL( activated() ) , this, SLOT( onRaytraceAction() ) );
+  myRaytraceActions.insert( ToolReflectionsId, a );
+  myRaytraceBar->addAction( a );
+
+  a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_ANTIALIASING") ), QObject::tr("MNU_TOOL_ANTIALIASING"), this );
+  a->setToolTip( QObject::tr("TBR_TOOL_ANTIALIASING") );
+  a->setStatusTip( QObject::tr("TBR_TOOL_ANTIALIASING") );
+  a->setCheckable( true );
+  a->setChecked( false );
+  connect( a, SIGNAL( activated() ) , this, SLOT( onRaytraceAction() ) );
+  myRaytraceActions.insert( ToolAntialiasingId, a );
+  myRaytraceBar->addAction( a );
+
+#endif
+
   QSignalMapper* sm = new QSignalMapper( this );
   connect( sm, SIGNAL( mapped( int ) ), this, SLOT( onSetMaterial( int ) ) );
 
@@ -275,12 +326,22 @@ void ApplicationCommonWindow::windowsMenuAboutToShow()
 
   QString dir = getResourceDir() + QString( "/" );
 
-  a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_NEW3D" ) ), QObject::tr( "MNU_WINDOW_NEW3D" ), this );
+  a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_NEW3D" ) ), QObject::tr( "MNU_WINDOW_NEW3D_GL" ), this );
-  a->setToolTip( QObject::tr( "TBR_WINDOW_NEW3D" ) );
+  a->setToolTip( QObject::tr( "TBR_WINDOW_NEW3D_GL" ) );
-  a->setStatusTip( QObject::tr( "TBR_WINDOW_NEW3D" ) );
+  a->setStatusTip( QObject::tr( "TBR_WINDOW_NEW3D_GL" ) );
   connect( a, SIGNAL( triggered() ), this, SLOT( onCreateNewView() ) );
   myWindowPopup->addAction( a );
 
+#ifdef HAVE_OPENCL
+
+  a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_NEW3D" ) ), QObject::tr( "MNU_WINDOW_NEW3D_RT" ), this );
+  a->setToolTip( QObject::tr( "TBR_WINDOW_NEW3D_RT" ) );
+  a->setStatusTip( QObject::tr( "TBR_WINDOW_NEW3D_RT" ) );
+  connect( a, SIGNAL( activated() ), this, SLOT( onCreateNewViewRT() ) );
+  myWindowPopup->addAction( a );
+
+#endif
+
   a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_CASCADE" ) ), QObject::tr( "MNU_WINDOW_CASCADE" ), this );
   a->setToolTip( QObject::tr( "TBR_WINDOW_CASCADE" ) );
   a->setStatusTip( QObject::tr( "TBR_WINDOW_CASCADE" ) );
@@ -401,7 +462,25 @@ DocumentCommon* ApplicationCommonWindow::onNewDoc()
     return aDoc;
 }
 
-void ApplicationCommonWindow::onCloseWindow(){
+DocumentCommon* ApplicationCommonWindow::onNewDocRT()
+{
+    updateFileActions();
+    DocumentCommon* aDoc = createNewDocument();
+    aDoc->onCreateNewView(true);
+    onSelectionChanged();
+    connect( aDoc, SIGNAL( sendCloseDocument( DocumentCommon* ) ),
+             this, SLOT( onCloseDocument( DocumentCommon* ) ) );
+    connect( stWs, SIGNAL( windowActivated( QWidget* ) ),
+             this, SLOT( onWindowActivated( QWidget* ) ) );
+    connect( aDoc, SIGNAL( selectionChanged() ),
+             this, SLOT( onSelectionChanged() ) );
+    myDocuments.append( aDoc );
+    myStdActions.at( FileCloseId )->setEnabled( myDocuments.count() > 0 );
+    return aDoc;
+}
+
+void ApplicationCommonWindow::onCloseWindow()
+{
     MDIWindow* m = (MDIWindow*)stWs->activeWindow();
     if ( m )
     {
@@ -410,6 +489,27 @@ void ApplicationCommonWindow::onCloseWindow(){
     }
 }
 
+void ApplicationCommonWindow::onUseVBO()
+{
+    MDIWindow* w = ( MDIWindow* ) stWs->activeWindow();
+    
+    if ( NULL == w )
+      return;
+
+    Handle(AIS_InteractiveContext) aContextAIS = w->getDocument()->getContext();
+
+    if (aContextAIS.IsNull())
+      return;
+
+    Handle(OpenGl_GraphicDriver) aDriver =
+      Handle(OpenGl_GraphicDriver)::DownCast (aContextAIS->CurrentViewer()->Driver());
+
+    if (!aDriver.IsNull())
+    {
+      aDriver->ChangeOptions().vboDisable = Standard_True;
+    }
+}
+
 void ApplicationCommonWindow::onCloseDocument(DocumentCommon* theDoc)
 {
   myDocuments.removeAll( theDoc );
@@ -449,14 +549,38 @@ void ApplicationCommonWindow::onAbout()
 
 void ApplicationCommonWindow::onCreateNewView()
 {
-    DocumentCommon* doc = ((MDIWindow*) stWs->activeWindow())->getDocument();
+  MDIWindow* window = qobject_cast< MDIWindow* >( stWs->activeWindow() );
-    doc->onCreateNewView();
+  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 )
+  if (w == NULL)
-        ((MDIWindow*) w)->onWindowActivated();
+  {
+    return;
+  }
+  
+  MDIWindow* window = qobject_cast< MDIWindow* >(w);
+
+  window->onWindowActivated();
+
+#ifdef HAVE_OPENCL
+
+  myRaytraceActions.at( ToolShadowsId )->setChecked (window->ShadowsEnabled());
+  myRaytraceActions.at( ToolReflectionsId )->setChecked (window->ReflectionsEnabled());
+  myRaytraceActions.at( ToolAntialiasingId )->setChecked (window->AntialiasingEnabled());
+
+#endif
 }
 
 void ApplicationCommonWindow::onToolAction()
@@ -484,6 +608,36 @@ void ApplicationCommonWindow::onToolAction()
         doc->onDelete();
 }
 
+#ifdef HAVE_OPENCL
+
+void ApplicationCommonWindow::onRaytraceAction()
+{
+  QAction* sentBy = (QAction*) sender();
+  
+  DocumentCommon* doc = qobject_cast< MDIWindow* >(
+    ApplicationCommonWindow::getWorkspace()->activeWindow())->getDocument();
+
+  if( sentBy == myRaytraceActions.at( ToolShadowsId ) )
+  {
+    bool flag = myRaytraceActions.at( ToolShadowsId )->isChecked(); 
+    doc->onShadows( flag );
+  }
+
+  if( sentBy == myRaytraceActions.at( ToolReflectionsId ) )
+  {
+    bool flag = myRaytraceActions.at( ToolReflectionsId )->isChecked();
+    doc->onReflections( flag );
+  }
+
+  if( sentBy == myRaytraceActions.at( ToolAntialiasingId ) )
+  {
+    bool flag = myRaytraceActions.at( ToolAntialiasingId )->isChecked();
+    doc->onAntialiasing( flag );
+  }
+}
+
+#endif
+
 void ApplicationCommonWindow::onSelectionChanged()
 {
   QWorkspace* ws = ApplicationCommonWindow::getWorkspace();

samples/qt/Common/src/ApplicationCommon.h

@@ -10,13 +10,15 @@
 #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();
@@ -36,14 +38,22 @@ protected:
 public slots:
 	
   DocumentCommon*                 onNewDoc();
+  DocumentCommon*                 onNewDocRT();
   void                            onCloseWindow();
+  void                            onUseVBO();
 	virtual void                    onCloseDocument( DocumentCommon* theDoc );
   virtual void                    onSelectionChanged();
   virtual void                    onAbout();
   void                            onViewToolBar();
 	void                            onViewStatusBar();
   void                            onToolAction();
+#ifdef HAVE_OPENCL
+  void                            onRaytraceAction();
+#endif
 	void                            onCreateNewView();
+#ifdef HAVE_OPENCL
+  void                            onCreateNewViewRT();
+#endif
   void                            onWindowActivated ( QWidget * w );
 	void                            windowsMenuAboutToShow();
   void                            windowsMenuActivated( bool checked/*int id*/ );
@@ -67,11 +77,13 @@ private:
 
 	QList<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;

samples/qt/Common/src/Common-icon.ts

@@ -86,6 +86,14 @@
         <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>
@@ -126,5 +134,17 @@
         <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>

samples/qt/Common/src/Common-string.ts

@@ -94,6 +94,14 @@
         <source>MNU_FILE</source>
         <translation>&amp;File</translation>
     </message>
+    <message>
+        <source>MNU_PREFERENCES</source>
+        <translation>&amp;Preferences</translation>
+    </message>
+    <message>
+        <source>MNU_USE_VBO</source>
+        <translation>&amp;Use VBO</translation>
+    </message>
     <message>
         <source>MNU_GOLD</source>
         <translation>&amp;Gold</translation>
@@ -130,6 +138,18 @@
         <source>MNU_TOOL_TRANS</source>
         <translation>&amp;Transpatency</translation>
     </message>
+    <message>
+        <source>MNU_TOOL_SHADOWS</source>
+        <translation>&amp;Shadows</translation>
+    </message>
+    <message>
+        <source>MNU_TOOL_REFLECTIONS</source>
+        <translation>&amp;Reflections</translation>
+    </message>
+    <message>
+        <source>MNU_TOOL_ANTIALIASING</source>
+        <translation>&amp;Antialiasing</translation>
+    </message>
     <message>
         <source>BTN_BRASS</source>
         <translation>Brass</translation>
@@ -234,6 +254,10 @@
         <source>MNU_CH_BACK</source>
         <translation>&amp;Change Background</translation>
     </message>
+    <message>
+        <source>MNU_CH_ENV_MAP</source>
+        <translation>&amp;Environment Map</translation>
+    </message>
     <message>
         <source>TBR_CH_BACK</source>
         <translation>Change Background</translation>
@@ -316,11 +340,23 @@
     </message>
     <message>
         <source>TBR_WINDOW_NEW3D</source>
-        <translation>New 3d View</translation>
+        <translation>New 3D View</translation>
     </message>
     <message>
-        <source>MNU_WINDOW_NEW3D</source>
+        <source>TBR_WINDOW_NEW3D_GL</source>
-        <translation>&amp;New 3d View</translation>
+        <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>&amp;New RT 3D View</translation>
+    </message>
+    <message>
+        <source>MNU_WINDOW_NEW3D_RT</source>
+        <translation>&amp;New RT 3D View</translation>
     </message>
     <message>
         <source>MNU_STATUS_BAR</source>

samples/qt/Common/src/DocumentCommon.cxx

@@ -50,12 +50,10 @@ myNbViews( 0 )
   myViewer = Viewer( getenv("DISPLAY"), a3DName.ToExtString(), "", 1000.0,
 		               V3d_XposYnegZpos, Standard_True, Standard_True );
 
-  myViewer->Init();
 	myViewer->SetDefaultLights();
 	myViewer->SetLightOn();
 
 	myContext =new AIS_InteractiveContext(myViewer);
-	//onCreateNewView();
 }
 
 DocumentCommon::~DocumentCommon()
@@ -67,15 +65,16 @@ ApplicationCommonWindow* DocumentCommon::getApplication()
 	return myApp;
 }
 
-MDIWindow* DocumentCommon::createNewMDIWindow()
+MDIWindow* DocumentCommon::createNewMDIWindow( bool theRT )
 {
   QWorkspace* ws = myApp->getWorkspace();
-  return new MDIWindow( this, ws, 0 );
+  return new MDIWindow( this, ws, 0, theRT );
 }
-void DocumentCommon::onCreateNewView() 
+
+void DocumentCommon::onCreateNewView( bool theRT ) 
 {
 	QWorkspace* ws = myApp->getWorkspace();
-  MDIWindow* w = createNewMDIWindow();
+  MDIWindow* w = createNewMDIWindow( theRT );
 	
 	if( !w )
 	  return;
@@ -228,4 +227,48 @@ void DocumentCommon::onDelete()
     getApplication()->onSelectionChanged();
 }
 
+#ifdef HAVE_OPENCL
 
+void DocumentCommon::onShadows( int state )
+{
+  QWorkspace* ws = ApplicationCommonWindow::getWorkspace();
+
+  MDIWindow* window = qobject_cast< MDIWindow* >( ws->activeWindow() );
+
+  if( window == NULL )
+    return;
+
+  window->setRaytracedShadows( state );
+
+  myContext->UpdateCurrentViewer();
+}
+
+void DocumentCommon::onReflections( int state )
+{
+  QWorkspace* ws = ApplicationCommonWindow::getWorkspace();
+
+  MDIWindow* window = qobject_cast< MDIWindow* >( ws->activeWindow() );
+
+  if( window == NULL )
+    return;
+
+  window->setRaytracedReflections( state );
+
+  myContext->UpdateCurrentViewer();
+}
+
+void DocumentCommon::onAntialiasing( int state )
+{
+  QWorkspace* ws = ApplicationCommonWindow::getWorkspace();
+
+  MDIWindow* window = qobject_cast< MDIWindow* >( ws->activeWindow() );
+
+  if( window == NULL )
+    return;
+
+  window->setRaytracedAntialiasing( state );
+
+  myContext->UpdateCurrentViewer();
+}
+
+#endif

samples/qt/Common/src/DocumentCommon.h

@@ -28,7 +28,7 @@ public:
 	void                           fitAll();
 	
 protected:
-  virtual MDIWindow*                   createNewMDIWindow();
+  virtual MDIWindow*                   createNewMDIWindow( bool theRT = false );
 
 signals:
   void                           selectionChanged();
@@ -36,10 +36,17 @@ signals:
 
 public slots:
 	virtual void                   onCloseView( MDIWindow* );
-	virtual void                   onCreateNewView();
+	virtual void                   onCreateNewView( bool theRT = false );
 	virtual void                   onMaterial();
   virtual void                   onMaterial( int );
 	virtual void                   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();

samples/qt/Common/src/MDIWindow.cxx

@@ -24,9 +24,17 @@ MDIWindow::MDIWindow(View* aView,
 
   myView = aView;
 	myDocument = aDocument;
+
+#ifdef HAVE_OPENCL
+
+  myShadowsEnabled = true;
+  myReflectionsEnabled = true;
+  myAntialiasingEnabled = false;
+
+#endif
 }  
   
-MDIWindow::MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags )
+MDIWindow::MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags, bool theRT )
 : QMainWindow( parent, wflags )
 {
 	QFrame *vb = new QFrame( this );
@@ -39,7 +47,7 @@ MDIWindow::MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlag
 	setCentralWidget( vb );
 
 	myDocument = aDocument;
-	myView = new View( myDocument->getContext(), vb );
+	myView = new View( myDocument->getContext(), vb, theRT );
 	layout->addWidget( myView );
 
   connect( myView, SIGNAL( selectionChanged() ),
@@ -49,6 +57,14 @@ MDIWindow::MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlag
   resize( sizeHint() );
 
   setFocusPolicy( Qt::StrongFocus );
+
+#ifdef HAVE_OPENCL
+
+  myShadowsEnabled = true;
+  myReflectionsEnabled = true;
+  myAntialiasingEnabled = false;
+
+#endif
 }
 
 MDIWindow::~MDIWindow()
@@ -137,4 +153,26 @@ QSize MDIWindow::sizeHint() const
     return QSize( 450, 300 );
 }
 
+#ifdef HAVE_OPENCL
+
+void MDIWindow::setRaytracedShadows( int state )
+{
+  myView->setRaytracedShadows( state );
+  myShadowsEnabled = state;
+}
+
+void MDIWindow::setRaytracedReflections( int state )
+{
+  myView->setRaytracedReflections( state );
+  myReflectionsEnabled = state;
+}
+
+void MDIWindow::setRaytracedAntialiasing( int state )
+{
+  myView->setRaytracedAntialiasing( state );
+  myAntialiasingEnabled = state;
+}
+
+#endif
+
 

samples/qt/Common/src/MDIWindow.h

@@ -12,7 +12,7 @@ class COMMONSAMPLE_EXPORT MDIWindow: public QMainWindow
     Q_OBJECT
 
 public:
-  MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags );
+  MDIWindow( DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags, bool theRT = false );
   MDIWindow( View* aView, DocumentCommon* aDocument, QWidget* parent, Qt::WindowFlags wflags );
   ~MDIWindow();
 
@@ -20,6 +20,18 @@ public:
 	void                       fitAll();
   virtual QSize              sizeHint() const;
 
+#ifdef HAVE_OPENCL
+
+  void                       setRaytracedShadows( int state );
+  void                       setRaytracedReflections( int state );
+  void                       setRaytracedAntialiasing( int state );
+
+  bool                       ShadowsEnabled() { return myShadowsEnabled; }
+  bool                       ReflectionsEnabled() { return myReflectionsEnabled; }
+  bool                       AntialiasingEnabled() { return myAntialiasingEnabled; }
+
+#endif
+
 signals:
   void                       selectionChanged();
   void                       message(const QString&, int );
@@ -36,6 +48,14 @@ protected:
 protected:
   DocumentCommon*            myDocument;
   View*                      myView;
+
+#ifdef HAVE_OPENCL
+
+  bool                       myShadowsEnabled;
+  bool                       myReflectionsEnabled;
+  bool                       myAntialiasingEnabled;
+
+#endif
 };
 
 #endif

samples/qt/Common/src/View.cxx

@@ -11,12 +11,14 @@
 #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__)
@@ -52,9 +54,11 @@ static QCursor* globPanCursor = NULL;
 static QCursor* zoomCursor    = NULL;
 static QCursor* rotCursor     = NULL;
 
-View::View( Handle(AIS_InteractiveContext) theContext, QWidget* parent )
+View::View( Handle(AIS_InteractiveContext) theContext, QWidget* parent, bool theRT )
 : QWidget( parent ),
-myViewActions( 0 )
+myIsRT( theRT ),
+myViewActions( 0 ),
+myBackMenu( NULL )
 {
 #if !defined(_WIN32) && !defined(__WIN32__) && (!defined(__APPLE__) || defined(MACOSX_USE_GLX))
   //XSynchronize( x11Display(),true ); // it is possible to use QApplication::syncX();
@@ -63,6 +67,9 @@ myViewActions( 0 )
     myFirst = true;
     myContext = theContext;
 
+    //if (theRT)
+    //  myContext->SetDisplayMode( AIS_DisplayMode::AIS_Shaded, 1 );
+
     myXmin = 0;
     myYmin = 0;
     myXmax = 0;
@@ -160,6 +167,7 @@ myViewActions( 0 )
 
 View::~View()
 {
+  delete myBackMenu;
 }
 
 void View::init()
@@ -184,6 +192,9 @@ void View::init()
   }
   myView->SetBackgroundColor (Quantity_NOC_BLACK);
   myView->MustBeResized();
+
+  if (myIsRT)
+    myView->SetRaytracingMode();
 }
 
 void View::paintEvent( QPaintEvent *  )
@@ -299,6 +310,30 @@ void View::hlrOn()
     QApplication::restoreOverrideCursor();
 }
 
+void View::setRaytracedShadows( int state )
+{
+  if ( state )
+    myView->EnableRaytracedShadows();
+  else
+    myView->DisableRaytracedShadows();
+}
+
+void View::setRaytracedReflections( int state )
+{
+  if ( state )
+    myView->EnableRaytracedReflections();
+  else
+    myView->DisableRaytracedReflections();
+}
+
+void View::setRaytracedAntialiasing( int state )
+{
+  if ( state )
+    myView->EnableRaytracedAntialiasing();
+  else
+    myView->DisableRaytracedAntialiasing();
+}
+
 void View::updateToggled( bool isOn )
 {
     QAction* sentBy = (QAction*)sender();
@@ -847,14 +882,26 @@ void View::Popup( const int /*x*/, const int /*y*/ )
   }
   else
   {
-    QMenu* myBackMenu = new QMenu( 0 );
+    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( triggered() ), this, SLOT( onBackground() ) );
+      connect( a, SIGNAL( activated() ), this, SLOT( onBackground() ) );
 		  myBackMenu->addAction( a );  
       addItemInPopup(myBackMenu);
+
+      a = new QAction( QObject::tr("MNU_CH_ENV_MAP"), this );
+		  a->setToolTip( QObject::tr("TBR_CH_ENV_MAP") );
+      connect( a, SIGNAL( activated() ), this, SLOT( onEnvironmentMap() ) );
+      a->setCheckable( true );
+      a->setChecked( false );
+		  myBackMenu->addAction( a );  
+      addItemInPopup(myBackMenu);
+    }
+
     myBackMenu->exec( QCursor::pos() );
-    delete myBackMenu;
   }
   if ( w )
     w->setFocus();
@@ -946,6 +993,26 @@ void View::onBackground()
     myView->Redraw();
 }
 
+void View::onEnvironmentMap()
+{
+  if (myBackMenu->actions().at(1)->isChecked())
+  {
+    QString fileName = QFileDialog::getOpenFileName(this, tr("Open File"), "",
+                           tr("All Image Files (*.bmp *.gif *.jpg *.jpeg *.png *.tga)"));
+
+    Handle(Graphic3d_TextureEnv) aTexture = new Graphic3d_TextureEnv( fileName.toAscii().data() );
+
+    myView->SetTextureEnv (aTexture);
+    myView->SetSurfaceDetail (V3d_TEX_ENVIRONMENT);
+  }
+  else
+  {
+    myView->SetSurfaceDetail (V3d_TEX_NONE);
+  }
+  
+  myView->Redraw();
+}
+
 bool View::dump(Standard_CString theFile)
 {
   myView->Redraw();

samples/qt/Common/src/View.h

@@ -26,7 +26,9 @@ public:
                       ViewAxoId, ViewRotationId, ViewResetId, ViewHlrOffId, ViewHlrOnId };
 
     View( Handle(AIS_InteractiveContext) theContext,
-          QWidget* parent);
+          QWidget* parent,
+          bool theRT = false );
+
     ~View();
 
     virtual void                  init();
@@ -35,6 +37,10 @@ public:
     void                          noActiveActions();
     bool                          isShadingMode();
 
+    void                          setRaytracedShadows( int state );
+    void                          setRaytracedReflections( int state );
+    void                          setRaytracedAntialiasing( int state );
+
     static QString                GetMessages( int type,TopAbs_ShapeEnum aSubShapeType,
                                                TopAbs_ShapeEnum aShapeType );
     static QString                GetShapeType( TopAbs_ShapeEnum aShapeType );
@@ -70,6 +76,7 @@ public slots:
     void                          hlrOff();
     void                          updateToggled( bool );
     void                          onBackground();
+    void                          onEnvironmentMap();
 
 protected:
     virtual void                  paintEvent( QPaintEvent* );
@@ -107,6 +114,7 @@ private:
                                                  const int MaxX, const int MaxY, const bool Draw );
 
 private:
+    bool                            myIsRT;
     bool                            myFirst;
     bool		                        myDrawRect;           // set when a rect is used for selection or magnify 
     Handle(V3d_View)                myView;
@@ -119,6 +127,7 @@ private:
     Quantity_Factor                 myCurZoom;
     Standard_Boolean                myHlrModeIsOn;
     QList<QAction*>*                myViewActions;
+    QMenu*                          myBackMenu;
     QRubberBand*                    myRectBand; //!< selection rectangle rubber band
 };
 

src/Graphic3d/Graphic3d_CView.hxx

@@ -94,7 +94,12 @@ public:
     Backfacing  (0),
     GDisplayCB  (NULL),
     GClientData (NULL),
-    ptrFBO (NULL)
+    ptrFBO (NULL),
+    WasRedrawnGL (0),
+    IsRaytracing (0),
+    IsShadowsEnabled (1),
+    IsReflectionsEnabled (1),
+    IsAntialiasingEnabled (0)
   {
     //
   }
@@ -131,6 +136,21 @@ public:
 
   void* ptrFBO;
 
+  //! Was the window redrawn by standard OpenGL?
+  mutable int WasRedrawnGL;
+
+  //! Enables/disables OpenCL-based ray-tracing.
+  int IsRaytracing;
+
+  //! Enables/disables ray-traced sharp shadows.
+  int IsShadowsEnabled;
+  
+  //! Enables/disables ray-traced reflections.
+  int IsReflectionsEnabled;
+  
+  //! Enables/disables ray-traced adaptive anti-aliasing.
+  int IsAntialiasingEnabled;
+
 };
 
 const Handle(Standard_Type)& TYPE(Graphic3d_CView);

src/OpenGl/EXTERNLIB

@@ -6,6 +6,7 @@ CSF_objc
 CSF_Appkit
 CSF_IOKit
 CSF_OpenGlLibs
+CSF_OPENCL
 CSF_AviLibs
 CSF_FREETYPE
 CSF_GL2PS

src/OpenGl/FILES

@@ -133,3 +133,11 @@ OpenGl_ShaderStates.cxx
 Handle_OpenGl_ShaderObject.hxx
 Handle_OpenGl_ShaderProgram.hxx
 Handle_OpenGl_ShaderManager.hxx
+OpenGl_Cl.hxx
+OpenGl_AABB.hxx
+OpenGl_AABB.cxx
+OpenGl_SceneGeometry.hxx
+OpenGl_SceneGeometry.cxx
+OpenGl_RaytraceTypes.hxx
+OpenGl_RaytraceSource.cxx
+OpenGl_Workspace_Raytrace.cxx

src/OpenGl/OpenGl_AABB.cxx

@@ -0,0 +1,137 @@
+// Created on: 2013-08-27
+// Created by: Denis BOGOLEPOV
+// Copyright (c) 2013 OPEN CASCADE SAS
+//
+// The content of this file is subject to the Open CASCADE Technology Public
+// License Version 6.5 (the "License"). You may not use the content of this file
+// except in compliance with the License. Please obtain a copy of the License
+// at http://www.opencascade.org and read it completely before using this file.
+//
+// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its
+// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France.
+//
+// The Original Code and all software distributed under the License is
+// distributed on an "AS IS" basis, without warranty of any kind, and the
+// Initial Developer hereby disclaims all such warranties, including without
+// limitation, any warranties of merchantability, fitness for a particular
+// purpose or non-infringement. Please see the License for the specific terms
+// and conditions governing the rights and limitations under the License.
+
+
+#ifdef HAVE_CONFIG_H
+  #include <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

src/OpenGl/OpenGl_AABB.hxx

@@ -0,0 +1,88 @@
+// Created on: 2013-08-27
+// Created by: Denis BOGOLEPOV
+// Copyright (c) 2013 OPEN CASCADE SAS
+//
+// The content of this file is subject to the Open CASCADE Technology Public
+// License Version 6.5 (the "License"). You may not use the content of this file
+// except in compliance with the License. Please obtain a copy of the License
+// at http://www.opencascade.org and read it completely before using this file.
+//
+// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its
+// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France.
+//
+// The Original Code and all software distributed under the License is
+// distributed on an "AS IS" basis, without warranty of any kind, and the
+// Initial Developer hereby disclaims all such warranties, including without
+// limitation, any warranties of merchantability, fitness for a particular
+// purpose or non-infringement. Please see the License for the specific terms
+// and conditions governing the rights and limitations under the License.
+
+
+#ifndef _OpenGl_AABB_Header
+#define _OpenGl_AABB_Header
+
+#include <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

src/OpenGl/OpenGl_Caps.cxx

@@ -35,7 +35,8 @@ OpenGl_Caps::OpenGl_Caps()
 #else
   contextDebug      (Standard_False),
 #endif
-  contextNoAccel    (Standard_False)
+  contextNoAccel    (Standard_False),
+  keepArrayData     (Standard_False)
 {
   //
 }

src/OpenGl/OpenGl_Caps.hxx

@@ -67,6 +67,13 @@ public: //! @name context creation parameters
    */
   Standard_Boolean contextNoAccel;
 
+  /**
+   * Disables freeing CPU memory after building VBOs.
+   *
+   * OFF by default.
+   */
+  Standard_Boolean keepArrayData;
+
 public: //! @name class methods
 
   //! Default constructor - initialize with most optimal values.

src/OpenGl/OpenGl_Cl.hxx

@@ -0,0 +1,33 @@
+// Created on: 2013-10-15
+// Created by: Denis BOGOLEPOV
+// Copyright (c) 2013 OPEN CASCADE SAS
+//
+// The content of this file is subject to the Open CASCADE Technology Public
+// License Version 6.5 (the "License"). You may not use the content of this file
+// except in compliance with the License. Please obtain a copy of the License
+// at http://www.opencascade.org and read it completely before using this file.
+//
+// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its
+// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France.
+//
+// The Original Code and all software distributed under the License is
+// distributed on an "AS IS" basis, without warranty of any kind, and the
+// Initial Developer hereby disclaims all such warranties, including without
+// limitation, any warranties of merchantability, fitness for a particular
+// purpose or non-infringement. Please see the License for the specific terms
+// and conditions governing the rights and limitations under the License.
+
+#ifndef _OpenGl_Cl_H__
+#define _OpenGl_Cl_H__
+
+// cl_gl.h includes OpenGL headers - make sure our stuff is included in right order
+#include <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__

src/OpenGl/OpenGl_GraphicDriver.cxx

@@ -17,6 +17,11 @@
 // purpose or non-infringement. Please see the License for the specific terms
 // and conditions governing the rights and limitations under the License.
 
+
+#ifdef HAVE_CONFIG_H
+  #include <config.h>
+#endif
+
 #include <OpenGl_GraphicDriver.hxx>
 
 #include <OpenGl_Context.hxx>
@@ -154,6 +159,34 @@ Standard_Boolean OpenGl_GraphicDriver::SetImmediateModeDrawToFront (const Graphi
   return Standard_False;
 }
 
+// =======================================================================
+// function : GetOpenClDeviceInfo
+// purpose  : Returns information about device used for computations
+// =======================================================================
+#ifndef HAVE_OPENCL
+
+Standard_Boolean OpenGl_GraphicDriver::GetOpenClDeviceInfo (const Graphic3d_CView&,
+  NCollection_DataMap<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  :

src/OpenGl/OpenGl_GraphicDriver.hxx

@@ -319,6 +319,12 @@ public:
   //! Method to setup UserDraw callback
   Standard_EXPORT OpenGl_UserDrawCallback_t& UserDrawCallback();
 
+public:
+  
+  //! Returns information about OpenCL device used for computations.
+  Standard_EXPORT Standard_Boolean GetOpenClDeviceInfo (const Graphic3d_CView& theCView,
+                      NCollection_DataMap<TCollection_AsciiString, TCollection_AsciiString>& theInfo);
+
 private:
 
   //! Method to retrieve valid GL context.

src/OpenGl/OpenGl_GraphicDriver_7.cxx

@@ -203,6 +203,16 @@ void OpenGl_GraphicDriver::Redraw (const Graphic3d_CView& ACView,
                                    const Standard_Integer /*width*/, 
                                    const Standard_Integer /*height*/)
 {
+  if (!myCaps->vboDisable && ACView.IsRaytracing)
+  {
+    if (ACView.WasRedrawnGL)
+    {
+      myDeviceLostFlag = Standard_True;
+    }
+
+    myCaps->keepArrayData = Standard_True;
+  }
+
   const OpenGl_CView *aCView = (const OpenGl_CView *)ACView.ptrView;
   if (aCView)
   {

src/OpenGl/OpenGl_Group.cxx

@@ -17,16 +17,25 @@
 // 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  :
 // =======================================================================
+#ifndef HAVE_OPENCL
 OpenGl_Group::OpenGl_Group()
+#else
+OpenGl_Group::OpenGl_Group (const OpenGl_Structure* theAncestorStructure)
+#endif
 : myAspectLine(NULL),
   myAspectFace(NULL),
   myAspectMarker(NULL),
@@ -34,6 +43,11 @@ OpenGl_Group::OpenGl_Group ()
   myFirst(NULL),
   myLast(NULL)
 {
+#ifdef HAVE_OPENCL
+  myAncestorStructure = theAncestorStructure;
+  myIsRaytracable = Standard_False;
+  myModificationState = 0; // initial state
+#endif
 }
 
 // =======================================================================
@@ -89,6 +103,18 @@ void OpenGl_Group::SetAspectFace (const CALL_DEF_CONTEXTFILLAREA& theAspect,
     anAspectFace->SetAspect (theAspect);
     AddElement (TelNil/*TelAspectFace*/, anAspectFace);
   }
+
+#ifdef HAVE_OPENCL
+  if (myIsRaytracable)
+  {
+    myModificationState++;
+
+    if (myAncestorStructure != NULL)
+    {
+      myAncestorStructure->UpdateStateWithAncestorStructures();
+    }
+  }
+#endif
 }
 
 // =======================================================================
@@ -141,15 +167,29 @@ void OpenGl_Group::SetAspectText (const CALL_DEF_CONTEXTTEXT& theAspect,
 // function : AddElement
 // purpose  :
 // =======================================================================
-void OpenGl_Group::AddElement (const TelType AType, OpenGl_Element *AElem )
+void OpenGl_Group::AddElement (const TelType theType, OpenGl_Element *theElem)
 {
-  OpenGl_ElementNode *node = new OpenGl_ElementNode();
+  OpenGl_ElementNode *aNode = new OpenGl_ElementNode();
 
-  node->type = AType;
+  aNode->type = theType;
-  node->elem = AElem;
+  aNode->elem = theElem;
-  node->next = NULL;
+  aNode->next = NULL;
-  (myLast? myLast->next : myFirst) = node;
+  (myLast? myLast->next : myFirst) = aNode;
-  myLast = node;
+  myLast = aNode;
+
+#ifdef HAVE_OPENCL
+  if (OpenGl_Raytrace::IsRaytracedElement (aNode))
+  {
+    myModificationState++;
+    myIsRaytracable = Standard_True;
+
+    if (myAncestorStructure != NULL)
+    {
+      myAncestorStructure->UpdateStateWithAncestorStructures();
+      myAncestorStructure->SetRaytracableWithAncestorStructures();
+    }
+  }
+#endif
 }
 
 // =======================================================================

src/OpenGl/OpenGl_Group.hxx

@@ -33,6 +33,7 @@
 #include <OpenGl_tsm.hxx>
 
 class OpenGl_Group;
+class OpenGl_Structure;
 
 typedef NCollection_List<const OpenGl_Group*     > OpenGl_ListOfGroup;
 
@@ -46,10 +47,13 @@ struct OpenGl_ElementNode
 
 class OpenGl_Group : public OpenGl_Element
 {
-
 public:
 
+#ifndef HAVE_OPENCL
   OpenGl_Group();
+#else
+  OpenGl_Group (const OpenGl_Structure* theAncestorStructure);
+#endif
 
   void SetAspectLine   (const CALL_DEF_CONTEXTLINE&     theAspect, const Standard_Boolean IsGlobal = Standard_True);
   void SetAspectFace   (const CALL_DEF_CONTEXTFILLAREA& theAspect, const Standard_Boolean IsGlobal = Standard_True);
@@ -61,6 +65,22 @@ public:
   virtual void Render  (const Handle(OpenGl_Workspace)& theWorkspace) const;
   virtual void Release (const Handle(OpenGl_Context)&   theGlCtx);
 
+  //! Returns first OpenGL element node of the group.
+  const OpenGl_ElementNode* FirstNode() const { return myFirst; }
+
+  //! Returns OpenGL face aspect.
+  const OpenGl_AspectFace* AspectFace() const { return myAspectFace; }
+
+#ifdef HAVE_OPENCL
+
+  //! Returns modification state for ray-tracing.
+  Standard_Size ModificationState() const { return myModificationState; }
+
+  //! Is the group ray-tracable (contains ray-tracable elements)?
+  Standard_Boolean IsRaytracable() const { return myIsRaytracable; }
+
+#endif
+
 protected:
 
   virtual ~OpenGl_Group();
@@ -75,6 +95,12 @@ protected:
   OpenGl_ElementNode*    myFirst;
   OpenGl_ElementNode*    myLast;
   
+#ifdef HAVE_OPENCL
+  const OpenGl_Structure*  myAncestorStructure;
+  Standard_Boolean         myIsRaytracable;
+  Standard_Size            myModificationState;
+#endif
+
 public:
 
   DEFINE_STANDARD_ALLOC

src/OpenGl/OpenGl_LayerList.cxx

@@ -19,6 +19,10 @@
 // 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>
@@ -152,6 +156,10 @@ void OpenGl_LayerList::AddStructure (const OpenGl_Structure *theStructure,
 
   aList.Add (theStructure, thePriority);
   myNbStructures++;
+
+  // Note: In ray-tracing mode we don't modify modification
+  // state here. It is redundant, because the possible changes
+  // will be handled in the loop for structures
 }
 
 //=======================================================================
@@ -173,6 +181,14 @@ void OpenGl_LayerList::RemoveStructure (const OpenGl_Structure *theStructure,
   if (aList.Remove (theStructure) >= 0)
   {
     myNbStructures--;
+
+#ifdef HAVE_OPENCL
+    if (theStructure->IsRaytracable())
+    {
+      myModificationState++;
+    }
+#endif
+
     return;
   }
   
@@ -188,6 +204,14 @@ void OpenGl_LayerList::RemoveStructure (const OpenGl_Structure *theStructure,
     if (aScanList.Remove (theStructure) >= 0)
     {
       myNbStructures--;
+
+#ifdef HAVE_OPENCL
+      if (theStructure->IsRaytracable())
+      {
+        myModificationState++;
+      }
+#endif
+
       return;
     }
   }

src/OpenGl/OpenGl_LayerList.hxx

@@ -81,6 +81,16 @@ class OpenGl_LayerList
   //! Render this element
   void Render (const Handle(OpenGl_Workspace) &theWorkspace) const;
 
+  //! Returns the set of OpenGL Z-layers.
+  const OpenGl_SequenceOfLayers& Layers() const { return myLayers; }
+  
+#ifdef HAVE_OPENCL
+
+  //! Returns structure modification state (for ray-tracing).
+  Standard_Size ModificationState() const { return myModificationState; }
+
+#endif
+
  private:
   
   //! Get default layer
@@ -94,6 +104,10 @@ class OpenGl_LayerList
   Standard_Integer        myNbPriorities;
   Standard_Integer        myNbStructures;
 
+#ifdef HAVE_OPENCL
+  mutable Standard_Size   myModificationState;
+#endif
+
  public:
   DEFINE_STANDARD_ALLOC
 };

src/OpenGl/OpenGl_PrimitiveArray.cxx

@@ -149,7 +149,11 @@ Standard_Boolean OpenGl_PrimitiveArray::BuildVBO (const Handle(OpenGl_Workspace)
     }
   }
 
+  if (!aGlCtx->caps->keepArrayData)
+  {
     clearMemoryOwn();
+  }
+  
   return Standard_True;
 }
 

src/OpenGl/OpenGl_PriorityList.hxx

@@ -58,6 +58,9 @@ class OpenGl_PriorityList
   //! or less). Returns Standard_False if the list can not be accepted.
   Standard_Boolean Append (const OpenGl_PriorityList& theOther);
 
+  //! Returns array of OpenGL structures.
+  const OpenGl_ArrayOfStructure& ArrayOfStructures() const { return myArray; }
+
  protected:
 
   OpenGl_ArrayOfStructure myArray;

src/OpenGl/OpenGl_RaytraceSource.cxx

@@ -0,0 +1,940 @@
+// Created on: 2013-10-16
+// Created by: Denis BOGOLEPOV
+// Copyright (c) 2013 OPEN CASCADE SAS
+//
+// The content of this file is subject to the Open CASCADE Technology Public
+// License Version 6.5 (the "License"). You may not use the content of this file
+// except in compliance with the License. Please obtain a copy of the License
+// at http://www.opencascade.org and read it completely before using this file.
+//
+// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its
+// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France.
+//
+// The Original Code and all software distributed under the License is
+// distributed on an "AS IS" basis, without warranty of any kind, and the
+// Initial Developer hereby disclaims all such warranties, including without
+// limitation, any warranties of merchantability, fitness for a particular
+// purpose or non-infringement. Please see the License for the specific terms
+// and conditions governing the rights and limitations under the License.
+
+#ifdef HAVE_CONFIG_H
+  #include <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

src/OpenGl/OpenGl_RaytraceTypes.hxx

@@ -0,0 +1,46 @@
+// Created on: 2013-10-15
+// Created by: Denis BOGOLEPOV
+// Copyright (c) 2013 OPEN CASCADE SAS
+//
+// The content of this file is subject to the Open CASCADE Technology Public
+// License Version 6.5 (the "License"). You may not use the content of this file
+// except in compliance with the License. Please obtain a copy of the License
+// at http://www.opencascade.org and read it completely before using this file.
+//
+// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its
+// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France.
+//
+// The Original Code and all software distributed under the License is
+// distributed on an "AS IS" basis, without warranty of any kind, and the
+// Initial Developer hereby disclaims all such warranties, including without
+// limitation, any warranties of merchantability, fitness for a particular
+// purpose or non-infringement. Please see the License for the specific terms
+// and conditions governing the rights and limitations under the License.
+
+
+#ifndef _OpenGl_RaytraceTypes_Header
+#define _OpenGl_RaytraceTypes_Header
+
+#include <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

src/OpenGl/OpenGl_SceneGeometry.cxx

@@ -0,0 +1,761 @@
+// Created on: 2013-08-27
+// Created by: Denis BOGOLEPOV
+// Copyright (c) 2013 OPEN CASCADE SAS
+//
+// The content of this file is subject to the Open CASCADE Technology Public
+// License Version 6.5 (the "License"). You may not use the content of this file
+// except in compliance with the License. Please obtain a copy of the License
+// at http://www.opencascade.org and read it completely before using this file.
+//
+// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its
+// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France.
+//
+// The Original Code and all software distributed under the License is
+// distributed on an "AS IS" basis, without warranty of any kind, and the
+// Initial Developer hereby disclaims all such warranties, including without
+// limitation, any warranties of merchantability, fitness for a particular
+// purpose or non-infringement. Please see the License for the specific terms
+// and conditions governing the rights and limitations under the License.
+
+#ifdef HAVE_CONFIG_H
+  #include <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

src/OpenGl/OpenGl_SceneGeometry.hxx

@@ -0,0 +1,360 @@
+// Created on: 2013-08-27
+// Created by: Denis BOGOLEPOV
+// Copyright (c) 2013 OPEN CASCADE SAS
+//
+// The content of this file is subject to the Open CASCADE Technology Public
+// License Version 6.5 (the "License"). You may not use the content of this file
+// except in compliance with the License. Please obtain a copy of the License
+// at http://www.opencascade.org and read it completely before using this file.
+//
+// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its
+// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France.
+//
+// The Original Code and all software distributed under the License is
+// distributed on an "AS IS" basis, without warranty of any kind, and the
+// Initial Developer hereby disclaims all such warranties, including without
+// limitation, any warranties of merchantability, fitness for a particular
+// purpose or non-infringement. Please see the License for the specific terms
+// and conditions governing the rights and limitations under the License.
+
+#ifndef _OpenGl_SceneGeometry_Header
+#define _OpenGl_SceneGeometry_Header
+
+#ifdef HAVE_OPENCL
+
+#include <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

src/OpenGl/OpenGl_Structure.cxx

@@ -17,6 +17,10 @@
 // purpose or non-infringement. Please see the License for the specific terms
 // and conditions governing the rights and limitations under the License.
 
+#ifdef HAVE_CONFIG_H
+  #include <config.h>
+#endif
+
 
 #include <OpenGl_CappingAlgo.hxx>
 #include <OpenGl_Context.hxx>
@@ -144,6 +148,10 @@ OpenGl_Structure::OpenGl_Structure ()
   myNamedStatus(0),
   myZLayer(0)
 {
+#if HAVE_OPENCL
+  myIsRaytracable = Standard_False;
+  myModificationState = 0;
+#endif
 }
 
 // =======================================================================
@@ -161,12 +169,21 @@ OpenGl_Structure::~OpenGl_Structure()
 // function : SetTransformation
 // purpose  :
 // =======================================================================
-void OpenGl_Structure::SetTransformation(const float *AMatrix)
+void OpenGl_Structure::SetTransformation (const float *theMatrix)
 {
   if (!myTransformation)
+  {
     myTransformation = new OpenGl_Matrix();
+  }
 
-  matcpy( myTransformation->mat, AMatrix );
+  matcpy (myTransformation->mat, theMatrix);
+
+#ifdef HAVE_OPENCL
+  if (myIsRaytracable)
+  {
+    UpdateStateWithAncestorStructures();
+  }
+#endif
 }
 
 // =======================================================================
@@ -208,6 +225,13 @@ void OpenGl_Structure::SetAspectFace (const CALL_DEF_CONTEXTFILLAREA& theAspect)
     myAspectFace = new OpenGl_AspectFace();
   }
   myAspectFace->SetAspect (theAspect);
+
+#ifdef HAVE_OPENCL
+  if (myIsRaytracable)
+  {
+    UpdateStateWithAncestorStructures();
+  }
+#endif
 }
 
 // =======================================================================
@@ -249,7 +273,11 @@ void OpenGl_Structure::SetHighlightBox (const Handle(OpenGl_Context)& theGlCtx,
   }
   else
   {
+#ifndef HAVE_OPENCL
     myHighlightBox = new OpenGl_Group();
+#else
+    myHighlightBox = new OpenGl_Group (this);
+#endif
   }
 
   CALL_DEF_CONTEXTLINE aContextLine;
@@ -306,29 +334,151 @@ void OpenGl_Structure::ClearHighlightColor (const Handle(OpenGl_Context)& theGlC
   myHighlightColor = NULL;
 }
 
+// =======================================================================
+// function : SetNamedStatus
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::SetNamedStatus (const Standard_Integer aStatus)
+{
+  myNamedStatus = aStatus;
+
+#ifdef HAVE_OPENCL
+  if (myIsRaytracable)
+  {
+    UpdateStateWithAncestorStructures();
+  }
+#endif
+}
+
+#ifdef HAVE_OPENCL
+
+// =======================================================================
+// function : RegisterAncestorStructure
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::RegisterAncestorStructure (const OpenGl_Structure* theStructure) const
+{
+  for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+  {
+    if (anIt.Value() == theStructure)
+    {
+      return;
+    }    
+  }
+
+  myAncestorStructures.Append (theStructure);
+}
+
+// =======================================================================
+// function : UnregisterAncestorStructure
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::UnregisterAncestorStructure (const OpenGl_Structure* theStructure) const
+{
+  for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+  {
+    if (anIt.Value() == theStructure)
+    {
+      myAncestorStructures.Remove (anIt);
+      return;
+    }    
+  }
+}
+
+// =======================================================================
+// function : UpdateStateWithAncestorStructures
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::UpdateStateWithAncestorStructures() const
+{
+  myModificationState++;
+
+  for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+  {
+    anIt.Value()->UpdateStateWithAncestorStructures();
+  }
+}
+
+// =======================================================================
+// function : UpdateRaytracableWithAncestorStructures
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::UpdateRaytracableWithAncestorStructures() const
+{
+  myIsRaytracable = OpenGl_Raytrace::IsRaytracedStructure (this);
+
+  if (!myIsRaytracable)
+  {
+    for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+    {
+      anIt.Value()->UpdateRaytracableWithAncestorStructures();
+    }
+  }
+}
+
+// =======================================================================
+// function : SetRaytracableWithAncestorStructures
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::SetRaytracableWithAncestorStructures() const
+{
+  myIsRaytracable = Standard_True;
+
+  for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+  {
+    if (!anIt.Value()->IsRaytracable())
+    {
+      anIt.Value()->SetRaytracableWithAncestorStructures();
+    }
+  }
+}
+
+#endif
+
 // =======================================================================
 // function : Connect
 // purpose  :
 // =======================================================================
-void OpenGl_Structure::Connect (const OpenGl_Structure *AStructure)
+void OpenGl_Structure::Connect (const OpenGl_Structure *theStructure)
 {
-  Disconnect (AStructure);
+  Disconnect (theStructure);
-  myConnected.Append(AStructure);
+  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);
   while (its.More())
   {
     // Check for the given structure
-    if (its.Value() == AStructure)
+    if (its.Value() == theStructure)
     {
       myConnected.Remove (its);
+
+#ifdef HAVE_OPENCL
+      if (theStructure->IsRaytracable())
+      {
+        UpdateStateWithAncestorStructures();
+        UpdateRaytracableWithAncestorStructures();
+      }
+
+      theStructure->UnregisterAncestorStructure (this);
+#endif
+
       return;
     }
     its.Next();
@@ -342,7 +492,12 @@ void OpenGl_Structure::Disconnect (const OpenGl_Structure *AStructure)
 OpenGl_Group * OpenGl_Structure::AddGroup()
 {
   // Create new group
-  OpenGl_Group *g = new OpenGl_Group;
+#ifndef HAVE_OPENCL
+  OpenGl_Group *g = new OpenGl_Group();
+#else
+  OpenGl_Group *g = new OpenGl_Group (this);
+#endif
+
   myGroups.Append(g);
   return g;
 }
@@ -359,9 +514,18 @@ void OpenGl_Structure::RemoveGroup (const Handle(OpenGl_Context)& theGlCtx,
     // Check for the given group
     if (anIter.Value() == theGroup)
     {
-      // Delete object
-      OpenGl_Element::Destroy (theGlCtx, const_cast<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;
     }
   }
@@ -373,13 +537,29 @@ void OpenGl_Structure::RemoveGroup (const Handle(OpenGl_Context)& theGlCtx,
 // =======================================================================
 void OpenGl_Structure::Clear (const Handle(OpenGl_Context)& theGlCtx)
 {
+#ifdef HAVE_OPENCL
+  Standard_Boolean aRaytracableGroupDeleted (Standard_False);
+#endif
+
   // Release groups
   for (OpenGl_ListOfGroup::Iterator anIter (myGroups); anIter.More(); anIter.Next())
   {
+#ifdef HAVE_OPENCL
+    aRaytracableGroupDeleted |= anIter.Value()->IsRaytracable();
+#endif
+    
     // Delete objects
     OpenGl_Element::Destroy (theGlCtx, const_cast<OpenGl_Group*& > (anIter.ChangeValue()));
   }
   myGroups.Clear();
+
+#ifdef HAVE_OPENCL
+  if (aRaytracableGroupDeleted)
+  {
+    UpdateStateWithAncestorStructures();
+    UpdateRaytracableWithAncestorStructures();
+  }
+#endif
 }
 
 // =======================================================================

src/OpenGl/OpenGl_Structure.hxx

@@ -30,6 +30,7 @@
 
 #include <OpenGl_Group.hxx>
 #include <OpenGl_Matrix.hxx>
+#include <OpenGl_NamedStatus.hxx>
 
 #include <Graphic3d_SetOfHClipPlane.hxx>
 
@@ -39,6 +40,7 @@ typedef NCollection_List<const OpenGl_Structure* > OpenGl_ListOfStructure;
 
 class OpenGl_Structure : public OpenGl_Element
 {
+  friend class OpenGl_Group;
 
 public:
 
@@ -65,7 +67,9 @@ public:
 
   void ClearHighlightColor (const Handle(OpenGl_Context)& theGlCtx);
 
-  void SetNamedStatus (const Standard_Integer aStatus) { myNamedStatus = aStatus; }
+  void SetNamedStatus (const Standard_Integer aStatus);
+
+  Standard_Boolean IsVisible() const { return !(myNamedStatus & OPENGL_NS_HIDE); }
 
   void SetClipPlanes (const Graphic3d_SetOfHClipPlane& thePlanes) { myClipPlanes = thePlanes; }
 
@@ -87,17 +91,64 @@ public:
   virtual void Release (const Handle(OpenGl_Context)&   theGlCtx);
 
   //! This method releases GL resources without actual elements destruction.
-  //! As result structure could be correctly destroyed layter without GL context
+  //! As result structure could be correctly destroyed layer without GL context
   //! (after last window was closed for example).
   //!
   //! Notice however that reusage of this structure after calling this method is incorrect
   //! and will lead to broken visualization due to loosed data.
   void ReleaseGlResources (const Handle(OpenGl_Context)& theGlCtx);
 
+  //! Returns list of OpenGL groups.
+  const OpenGl_ListOfGroup& Groups() const { return myGroups; }
+
+  //! Returns list of connected OpenGL structures.
+  const OpenGl_ListOfStructure& ConnectedStructures() const { return myConnected; }
+
+  //! Returns OpenGL face aspect.
+  const OpenGl_AspectFace* AspectFace() const { return myAspectFace; }
+
+  //! Returns OpenGL transformation matrix.
+  const OpenGl_Matrix* Transformation() const { return myTransformation; }
+  
+  //! Returns OpenGL persistent translation.
+  const TEL_TRANSFORM_PERSISTENCE* PersistentTranslation() const { return myTransPers; }
+
+#ifdef HAVE_OPENCL
+
+  //! Returns structure modification state (for ray-tracing).
+  Standard_Size ModificationState() const { return myModificationState; }
+
+  //! Resets structure modification state (for ray-tracing)
+  void ResetModificationState() const { myModificationState = 0; }
+
+  //! Is the structure ray-tracable (contains ray-tracable elements)?
+  Standard_Boolean IsRaytracable() const { return myIsRaytracable; }
+
+#endif
+
 protected:
 
   virtual ~OpenGl_Structure();
 
+#ifdef HAVE_OPENCL
+
+  //! Registers ancestor connected structure (for updating ray-tracing state).
+  void RegisterAncestorStructure (const OpenGl_Structure* theStructure) const;
+
+  //! Unregisters ancestor connected structure (for updating ray-tracing state).
+  void UnregisterAncestorStructure (const OpenGl_Structure* theStructure) const;
+
+  //! Updates modification state for structure and its parents.
+  void UpdateStateWithAncestorStructures() const;
+
+  //! Updates ray-tracable status for structure and its parents.
+  void UpdateRaytracableWithAncestorStructures() const;
+
+  //! Sets ray-tracable status for structure and its parents.
+  void SetRaytracableWithAncestorStructures() const;
+
+#endif
+
 protected:
 
   //Structure_LABBegin
@@ -119,6 +170,12 @@ protected:
   OpenGl_ListOfGroup         myGroups;
   Graphic3d_SetOfHClipPlane  myClipPlanes;
 
+#ifdef HAVE_OPENCL
+  mutable OpenGl_ListOfStructure   myAncestorStructures;
+  mutable Standard_Boolean         myIsRaytracable;
+  mutable Standard_Size            myModificationState;
+#endif
+
 public:
 
   DEFINE_STANDARD_ALLOC

src/OpenGl/OpenGl_View.cxx

@@ -17,6 +17,11 @@
 // purpose or non-infringement. Please see the License for the specific terms
 // and conditions governing the rights and limitations under the License.
 
+
+#ifdef HAVE_CONFIG_H
+  #include <config.h>
+#endif
+
 #include <NCollection_Mat4.hxx>
 
 #include <OpenGl_Context.hxx>
@@ -123,6 +128,10 @@ OpenGl_View::OpenGl_View (const CALL_DEF_VIEWCONTEXT &AContext)
       myIntShadingMethod = TEL_SM_FLAT;
       break;
   }
+
+#ifdef HAVE_OPENCL
+  myModificationState = 1; // initial state
+#endif
 }
 
 /*----------------------------------------------------------------------*/
@@ -167,6 +176,19 @@ void OpenGl_View::SetTextureEnv (const Handle(OpenGl_Context)&       theCtx,
   Handle(Image_PixMap) anImage = theTexture->GetImage();
   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
 }
 
 /*----------------------------------------------------------------------*/

src/OpenGl/OpenGl_View.hxx

@@ -109,7 +109,7 @@ class OpenGl_View : public MMgt_TShared
 
   void SetTextureEnv (const Handle(OpenGl_Context)&       theCtx,
                       const Handle(Graphic3d_TextureEnv)& theTexture);
-  void SetSurfaceDetail (const Visual3d_TypeOfSurfaceDetail AMode) { mySurfaceDetail = AMode; }
+  void SetSurfaceDetail (const Visual3d_TypeOfSurfaceDetail AMode);
   void SetBackfacing (const Standard_Integer AMode);
   void SetLights (const CALL_DEF_VIEWCONTEXT &AContext);
   void SetAntiAliasing (const Standard_Boolean AMode) { myAntiAliasing = AMode; }
@@ -189,6 +189,26 @@ class OpenGl_View : public MMgt_TShared
                const Aspect_CLayer2d&               theCUnderLayer,
                const Aspect_CLayer2d&               theCOverLayer);
 
+
+  void DrawBackground (const Handle(OpenGl_Workspace)& theWorkspace);
+
+  //! Returns list of OpenGL Z-layers.
+  const OpenGl_LayerList& LayerList() const { return myZLayers; }
+
+  //! Returns list of openGL light sources.
+  const OpenGl_ListOfLight& LightList() const { return myLights; }
+
+  //! Returns OpenGL environment map.
+  const Handle(OpenGl_Texture)& TextureEnv() const { return myTextureEnv; }
+
+  //! Returns visualization mode for objects in the view.
+  Visual3d_TypeOfSurfaceDetail SurfaceDetail() const { return mySurfaceDetail; }
+
+#ifdef HAVE_OPENCL
+  //! Returns modification state for ray-tracing.
+  Standard_Size ModificationState() const { return myModificationState; }
+#endif
+
 public:
 
   DEFINE_STANDARD_RTTI(OpenGl_View) // Type definition
@@ -250,6 +270,10 @@ public:
   Standard_Boolean myViewMappingChanged;
   Standard_Boolean myLightSourcesChanged;
 
+#ifdef HAVE_OPENCL
+  Standard_Size myModificationState;
+#endif
+
  public:
   DEFINE_STANDARD_ALLOC
 };

src/OpenGl/OpenGl_View_2.cxx

@@ -636,67 +636,8 @@ call_util_mat_mul( matrix3 mat_a, matrix3 mat_b, matrix3 mat_c)
 
 /*----------------------------------------------------------------------*/
 
-//call_func_redraw_all_structs_proc
+void OpenGl_View::DrawBackground (const Handle(OpenGl_Workspace) &AWorkspace)
-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
 
@@ -891,6 +832,76 @@ void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext,
     else
       glDisable (GL_DITHER);
   }
+}
+
+/*----------------------------------------------------------------------*/
+
+//call_func_redraw_all_structs_proc
+void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext,
+                          const Handle(OpenGl_Workspace) &AWorkspace,
+                          const Graphic3d_CView& ACView,
+                          const Aspect_CLayer2d& ACUnderLayer,
+                          const Aspect_CLayer2d& ACOverLayer)
+{
+  // Store and disable current clipping planes
+  const Handle(OpenGl_Context)& aContext = AWorkspace->GetGlContext();
+  const Standard_Integer aMaxClipPlanes = aContext->MaxClipPlanes();
+  const GLenum lastid = GL_CLIP_PLANE0 + aMaxClipPlanes;
+  OPENGL_CLIP_PLANE *oldPlanes = new OPENGL_CLIP_PLANE[aMaxClipPlanes];
+  OPENGL_CLIP_PLANE *ptrPlane = oldPlanes;
+  GLenum planeid = GL_CLIP_PLANE0;
+  for ( ; planeid < lastid; planeid++, ptrPlane++ )
+  {
+    glGetClipPlane( planeid, ptrPlane->Equation );
+    if ( ptrPlane->isEnabled )
+    {
+      glDisable( planeid );
+      ptrPlane->isEnabled = GL_TRUE;
+    }
+    else
+    {
+      ptrPlane->isEnabled = GL_FALSE;
+    }
+  }
+
+  // Set OCCT state uniform variables
+
+  if (!aContext->ShaderManager()->IsEmpty())
+  {
+    if (myLightSourcesChanged)
+    {
+      aContext->ShaderManager()->UpdateLightSourceStateTo (&myLights);
+      myLightSourcesChanged = Standard_False;
+    }
+
+    if (myViewMappingChanged)
+    {
+      aContext->ShaderManager()->UpdateProjectionStateTo (myMappingMatrix);
+      myViewMappingChanged = Standard_False;
+    }
+
+    if (myOrientationChanged)
+    {
+      aContext->ShaderManager()->UpdateWorldViewStateTo (myOrientationMatrix);
+      myOrientationChanged = Standard_False;
+    }
+
+    if (aContext->ShaderManager()->ModelWorldState().Index() == 0)
+    {
+      Tmatrix3 aModelWorldState = { { 1.f, 0.f, 0.f, 0.f },
+                                    { 0.f, 1.f, 0.f, 0.f },
+                                    { 0.f, 0.f, 1.f, 0.f },
+                                    { 0.f, 0.f, 0.f, 1.f } };
+
+      aContext->ShaderManager()->UpdateModelWorldStateTo (aModelWorldState);
+    }
+  }
+
+  /////////////////////////////////////////////////////////////////////////////
+  // Step 1: Prepare for redraw
+
+  // Render background
+  DrawBackground (AWorkspace);
 
   // Switch off lighting by default
   glDisable(GL_LIGHTING);
@@ -1028,8 +1039,6 @@ D = -[Px,Py,Pz] dot |Nx|
 
   // Apply clipping planes
   {
-    const Handle(OpenGl_Context)& aContext = AWorkspace->GetGlContext();
-
     if (myZClip.Back.IsOn || myZClip.Front.IsOn)
     {
       const GLdouble ramp = myExtra.map.fpd - myExtra.map.bpd;

src/OpenGl/OpenGl_Workspace.cxx

@@ -17,6 +17,11 @@
 // purpose or non-infringement. Please see the License for the specific terms
 // and conditions governing the rights and limitations under the License.
 
+
+#ifdef HAVE_CONFIG_H
+  #include <config.h>
+#endif
+
 #include <OpenGl_GlCore15.hxx>
 
 #include <InterfaceGraphic.hxx>
@@ -28,6 +33,7 @@
 #include <OpenGl_Context.hxx>
 #include <OpenGl_FrameBuffer.hxx>
 #include <OpenGl_Texture.hxx>
+#include <OpenGl_View.hxx>
 #include <OpenGl_Workspace.hxx>
 #include <OpenGl_Element.hxx>
 
@@ -116,6 +122,21 @@ OpenGl_Workspace::OpenGl_Workspace (const Handle(OpenGl_Display)& theDisplay,
 
   // Polygon Offset
   EnablePolygonOffset();
+
+#ifdef HAVE_OPENCL
+
+  myComputeInitStatus = OpenGl_CLIS_NONE;
+
+  myViewModificationStatus = 0;
+  myLayersModificationStatus = 0;
+
+  myRaytraceOutputTexture[0] = 0;
+  myRaytraceOutputTexture[1] = 0;
+
+  myIsRaytraceDataValid = Standard_False;
+  myToUpdateRaytraceData = Standard_False;
+
+#endif
 }
 
 // =======================================================================
@@ -135,6 +156,9 @@ Standard_Boolean OpenGl_Workspace::SetImmediateModeDrawToFront (const Standard_B
 // =======================================================================
 OpenGl_Workspace::~OpenGl_Workspace()
 {
+#ifdef HAVE_OPENCL
+  ReleaseOpenCL();
+#endif
 }
 
 // =======================================================================
@@ -483,12 +507,30 @@ void OpenGl_Workspace::Redraw (const Graphic3d_CView& theCView,
     toSwap = 0; // no need to swap buffers
   }
 
+#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
+  {
+    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)
   {
     aFrameBuffer->UnbindBuffer (aGlCtx);

src/OpenGl/OpenGl_Workspace.hxx

@@ -20,6 +20,13 @@
 #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>
 
@@ -48,6 +55,9 @@
 #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>
 
@@ -62,8 +72,8 @@ class OpenGl_Structure;
 class OpenGl_Element;
 class Image_PixMap;
 
-//! Reprepsents window with GL context.
+//! Represents window with GL context.
-//! Provides methods to render primitives and maintan GL state.
+//! Provides methods to render primitives and maintain GL state.
 class OpenGl_Workspace : public OpenGl_Window
 {
 public:
@@ -208,6 +218,171 @@ protected:
   void setTextureParams (Handle(OpenGl_Texture)&                theTexture,
                          const Handle(Graphic3d_TextureParams)& theParams);
 
+#ifdef HAVE_OPENCL
+
+public:
+
+  //! Returns information about OpenCL device used for computations.
+  Standard_Boolean GetOpenClDeviceInfo (
+    NCollection_DataMap<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;

src/TKOpenGl/EXTERNLIB

@@ -5,6 +5,7 @@ CSF_OpenGlLibs
 CSF_objc
 CSF_Appkit
 CSF_IOKit
+CSF_OPENCL
 CSF_FREETYPE
 CSF_GL2PS
 CSF_user32

src/V3d/V3d_View.cdl

@@ -1580,6 +1580,38 @@ is
      ---Purpose: Get clip planes.
      -- @return sequence clip planes that have been set for the view
 
+        SetRaytracingMode (me : mutable) is static;
+     ---Level: Public
+     ---Purpose: enables OpenCL-based ray-tracing mode
+
+        SetRasterizationMode (me : mutable) is static;
+     ---Level: Public
+     ---Purpose: enables OpenGL-based rasterization mode
+
+        EnableRaytracedShadows (me : mutable) is static;
+     ---Level: Public
+     ---Purpose: enables sharp shadows in OpenCL-based ray-tracing mode
+
+        EnableRaytracedReflections (me : mutable) is static;
+     ---Level: Public
+     ---Purpose: enables specular reflections in OpenCL-based ray-tracing mode
+
+        EnableRaytracedAntialiasing (me : mutable) is static;
+     ---Level: Public
+     ---Purpose: enables antialiasing in OpenCL-based ray-tracing mode
+
+        DisableRaytracedShadows (me : mutable) is static;
+     ---Level: Public
+     ---Purpose: disables sharp shadows in OpenCL-based ray-tracing mode
+
+        DisableRaytracedReflections (me : mutable) is static;
+     ---Level: Public
+     ---Purpose: disables specular reflections in OpenCL-based ray-tracing mode
+
+        DisableRaytracedAntialiasing (me : mutable) is static;
+     ---Level: Public
+     ---Purpose: disables antialiasing in OpenCL-based ray-tracing mode
+
 fields
 
         MyType :                TypeOfView from V3d is protected ;

src/V3d/V3d_View_5.cxx

@@ -57,6 +57,62 @@ Standard_Boolean V3d_View::IsGLLightEnabled() const
   return MyView->IsGLLightEnabled();
 }
 
+void V3d_View::SetRaytracingMode()
+{
+  Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView();
+
+  cView->IsRaytracing = 1;
+}
+
+void V3d_View::SetRasterizationMode()
+{
+  Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView();
+
+  cView->IsRaytracing = 0;
+}
+
+void V3d_View::EnableRaytracedShadows()
+{
+  Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView();
+
+  cView->IsShadowsEnabled = 1;
+}
+
+void V3d_View::EnableRaytracedReflections()
+{
+  Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView();
+
+  cView->IsReflectionsEnabled = 1;
+}
+
+void V3d_View::EnableRaytracedAntialiasing()
+{
+  Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView();
+
+  cView->IsAntialiasingEnabled = 1;
+}
+
+void V3d_View::DisableRaytracedShadows()
+{
+  Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView();
+
+  cView->IsShadowsEnabled = 0;
+}
+
+void V3d_View::DisableRaytracedReflections()
+{
+  Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView();
+
+  cView->IsReflectionsEnabled = 0;
+}
+
+void V3d_View::DisableRaytracedAntialiasing()
+{
+  Graphic3d_CView* cView = (Graphic3d_CView* )MyView->CView();
+
+  cView->IsAntialiasingEnabled = 0;
+}
+
 void V3d_View::SetLayerMgr(const Handle(V3d_LayerMgr)& aMgr)
 {
   MyLayerMgr = aMgr;

src/ViewerTest/ViewerTest_ViewerCommands.cxx

@@ -5275,6 +5275,183 @@ static int VSetTextureMode (Draw_Interpretor& theDi, Standard_Integer theArgsNb,
   return 0;
 }
 
+//==============================================================================
+//function : VClInfo
+//purpose  : Prints info about active OpenCL device
+//==============================================================================
+
+static Standard_Integer VClInfo (Draw_Interpretor& theInterpretor,
+                                 Standard_Integer,
+                                 const char**)
+{
+#ifndef HAVE_OPENCL
+
+  theInterpretor << "OCCT was compiled without OpenCL support!\n";
+
+#else
+
+  Handle(AIS_InteractiveContext) aContextAIS = ViewerTest::GetAISContext();
+
+  if (aContextAIS.IsNull())
+  {
+    theInterpretor << "Call vinit before!\n";
+    return 1;
+  }
+
+  Handle(OpenGl_GraphicDriver) aDrv = Handle(OpenGl_GraphicDriver)::DownCast (aContextAIS->CurrentViewer()->Driver());
+  
+  Graphic3d_CView* aCView = (Graphic3d_CView*) ViewerTest::CurrentView()->View()->CView();
+
+  NCollection_DataMap<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  :
@@ -5548,4 +5725,13 @@ void ViewerTest::ViewerCommands(Draw_Interpretor& theCommands)
     "  2 - all textures enabled.\n"
     "  this command sets texture details mode for the specified view.\n"
     , __FILE__, VSetTextureMode, group);
+  theCommands.Add("vraytrace",
+    "vraytrace 0|1",
+    __FILE__,VRaytrace,group);
+  theCommands.Add("vclinfo",
+    "vclinfo",
+    __FILE__,VClInfo,group);
+  theCommands.Add("vsetraytracemode",
+    "vsetraytracemode [shad=0|1] [refl=0|1] [aa=0|1]",
+    __FILE__,VSetRaytraceMode,group);
 }

tests/bugs/parse.rules

@@ -1,4 +1,5 @@
 FAILED /\bFaulty\b/ bad shape
 IGNORE /^Error [23]d = [\d.-]+/ debug output of blend command
 SKIPPED /Error: unsupported locale specification/ locale is unavailable on tested system
+SKIPPED /OCCT was compiled without OpenCL support!/
 OK /Relative error of mass computation/ message from vprops

tests/bugs/vis/bug24130

@@ -0,0 +1,48 @@
+puts "TODO OCC24130 Debian60-64: OCCT was compiled without OpenCL support!"
+puts "TODO OCC24130 Windows: TKOpenGl | Type\: Error | ID\: 0 | Severity\: High | Message\:"
+
+puts "========"
+puts "OCC24130 Implementing ray tracing visualization core"
+puts "========"
+
+# custom shapes
+set aShape1 [locate_data_file occ/Top.brep]
+set aShape2 [locate_data_file occ/Bottom.brep]
+
+# setup 3D viewer content
+vinit name=View1 w=512 h=512
+vglinfo
+
+vvbo 0
+vsetdispmode 1
+vsetgradientbg 180 200 255 180 180 180 2
+restore $aShape1 s1
+restore $aShape2 s2
+vdisplay s1 s2
+vsetmaterial s1 Silver
+vsetmaterial s2 Pewter
+vfit
+
+# activate ray-tracing
+vraytrace 1
+vclinfo
+
+set aModeNum 0
+for { set aAAMode 0 } { $aAAMode <= 1 } { incr aAAMode } {
+  for { set aReflMode 0 } { $aReflMode <= 1 } { incr aReflMode } {
+    for { set aShadMode 0 } { $aShadMode <= 1 } { incr aShadMode } {
+      vsetraytracemode shad=$aShadMode refl=$aReflMode aa=$aAAMode
+      vdump $imagedir/${casename}_${aModeNum}.png
+      incr aModeNum
+    }
+  }
+}
+
+vtextureenv on 5
+for { set aAAMode 0 } { $aAAMode <= 1 } { incr aAAMode } {
+  for { set aShadMode 0 } { $aShadMode <= 1 } { incr aShadMode } {
+    vsetraytracemode shad=$aShadMode refl=1 aa=$aAAMode
+    vdump $imagedir/${casename}_${aModeNum}.png
+    incr aModeNum
+  }
+}