0024130: Implementing ray tracing visualization core
authordbp <dbp@opencascade.com>
Thu, 31 Oct 2013 11:35:18 +0000 (15:35 +0400)
committerbugmaster <bugmaster@opencascade.com>
Thu, 31 Oct 2013 14:02:12 +0000 (18:02 +0400)
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.

52 files changed:
.gitattributes
adm/UDLIST
samples/qt/Common/res/antialiasing.png [new file with mode: 0644]
samples/qt/Common/res/reflections.png [new file with mode: 0644]
samples/qt/Common/res/shadows.png [new file with mode: 0644]
samples/qt/Common/src/ApplicationCommon.cxx
samples/qt/Common/src/ApplicationCommon.h
samples/qt/Common/src/Common-icon.ts
samples/qt/Common/src/Common-string.ts
samples/qt/Common/src/DocumentCommon.cxx
samples/qt/Common/src/DocumentCommon.h
samples/qt/Common/src/MDIWindow.cxx
samples/qt/Common/src/MDIWindow.h
samples/qt/Common/src/View.cxx
samples/qt/Common/src/View.h
samples/qt/IESample/IESample-vc10.sln
src/Graphic3d/Graphic3d_CView.hxx
src/OpenGl/EXTERNLIB
src/OpenGl/FILES
src/OpenGl/OpenGl_AABB.cxx [new file with mode: 0644]
src/OpenGl/OpenGl_AABB.hxx [new file with mode: 0644]
src/OpenGl/OpenGl_Caps.cxx
src/OpenGl/OpenGl_Caps.hxx
src/OpenGl/OpenGl_Cl.hxx [new file with mode: 0644]
src/OpenGl/OpenGl_Display_2.cxx
src/OpenGl/OpenGl_GraphicDriver.cxx
src/OpenGl/OpenGl_GraphicDriver.hxx
src/OpenGl/OpenGl_GraphicDriver_7.cxx
src/OpenGl/OpenGl_Group.cxx
src/OpenGl/OpenGl_Group.hxx
src/OpenGl/OpenGl_LayerList.cxx
src/OpenGl/OpenGl_LayerList.hxx
src/OpenGl/OpenGl_PrimitiveArray.cxx
src/OpenGl/OpenGl_PriorityList.hxx
src/OpenGl/OpenGl_RaytraceSource.cxx [new file with mode: 0644]
src/OpenGl/OpenGl_RaytraceTypes.hxx [new file with mode: 0644]
src/OpenGl/OpenGl_SceneGeometry.cxx [new file with mode: 0644]
src/OpenGl/OpenGl_SceneGeometry.hxx [new file with mode: 0644]
src/OpenGl/OpenGl_Structure.cxx
src/OpenGl/OpenGl_Structure.hxx
src/OpenGl/OpenGl_View.cxx
src/OpenGl/OpenGl_View.hxx
src/OpenGl/OpenGl_View_2.cxx
src/OpenGl/OpenGl_Workspace.cxx
src/OpenGl/OpenGl_Workspace.hxx
src/OpenGl/OpenGl_Workspace_Raytrace.cxx [new file with mode: 0644]
src/TKOpenGl/EXTERNLIB
src/V3d/V3d_View.cdl
src/V3d/V3d_View_5.cxx
src/ViewerTest/ViewerTest_ViewerCommands.cxx
tests/bugs/parse.rules
tests/bugs/vis/bug24130 [new file with mode: 0644]

index e805ec9..6f32bdc 100644 (file)
@@ -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
index eb7824c..c6cea58 100644 (file)
@@ -488,3 +488,4 @@ p BOPAlgo
 p BOPDS
 p BOPCol
 p BOPInt
+r OpenCL
diff --git a/samples/qt/Common/res/antialiasing.png b/samples/qt/Common/res/antialiasing.png
new file mode 100644 (file)
index 0000000..da8e504
Binary files /dev/null and b/samples/qt/Common/res/antialiasing.png differ
diff --git a/samples/qt/Common/res/reflections.png b/samples/qt/Common/res/reflections.png
new file mode 100644 (file)
index 0000000..ad48a95
Binary files /dev/null and b/samples/qt/Common/res/reflections.png differ
diff --git a/samples/qt/Common/res/shadows.png b/samples/qt/Common/res/shadows.png
new file mode 100644 (file)
index 0000000..7562c12
Binary files /dev/null and b/samples/qt/Common/res/shadows.png differ
index 186363f..5751c69 100755 (executable)
@@ -12,6 +12,9 @@
 #include <QApplication>
 #include <QSignalMapper>
 
+#include <Graphic3d_GraphicDriver.hxx>
+#include <OpenGl_GraphicDriver.hxx>
+
 #include <stdlib.h>
 
 static ApplicationCommonWindow* stApp;
@@ -56,13 +59,12 @@ void ApplicationCommonWindow::createStandardOperations()
 
   QString dir = getResourceDir() + QString( "/" );
 
-  newIcon = QPixmap( dir+QObject::tr("ICON_NEW") );
-  helpIcon = QPixmap( dir+QObject::tr("ICON_HELP") );
-  closeIcon = QPixmap( dir+QObject::tr("ICON_CLOSE") );
+  newIcon = QPixmap( dir + QObject::tr("ICON_NEW") );
+  helpIcon = QPixmap( dir + QObject::tr("ICON_HELP") );
+  closeIcon = QPixmap( dir + QObject::tr("ICON_CLOSE") );
 
-  QAction * fileNewAction, * fileCloseAction, * fileQuitAction,
-          * viewToolAction, * viewStatusAction,
-                 * helpAboutAction;
+  QAction * fileNewAction, * fileCloseAction, * filePrefUseVBOAction,
+          * fileQuitAction, * viewToolAction, * viewStatusAction, * helpAboutAction;
 
   fileNewAction = new QAction( newIcon, QObject::tr("MNU_NEW"), this );
   fileNewAction->setToolTip( QObject::tr("TBR_NEW") );
@@ -78,6 +80,14 @@ void ApplicationCommonWindow::createStandardOperations()
   connect( fileCloseAction, SIGNAL( triggered() ) , this, SLOT( onCloseWindow() ) );
   myStdActions.insert( FileCloseId, fileCloseAction );
 
+  filePrefUseVBOAction = new QAction( QObject::tr("MNU_USE_VBO"), this );
+  filePrefUseVBOAction->setToolTip( QObject::tr("TBR_USE_VBO") );
+  filePrefUseVBOAction->setStatusTip( QObject::tr("TBR_USE_VBO") );
+  filePrefUseVBOAction->setCheckable( true );
+  filePrefUseVBOAction->setChecked( true );
+  connect( filePrefUseVBOAction, SIGNAL( activated() ) , this, SLOT( onUseVBO() ) );
+  myStdActions.insert( FilePrefUseVBOId, filePrefUseVBOAction );
+
   fileQuitAction = new QAction( QObject::tr("MNU_QUIT"), this );
   fileQuitAction->setToolTip( QObject::tr("TBR_QUIT") );
   fileQuitAction->setStatusTip( QObject::tr("TBR_QUIT") );
@@ -108,12 +118,18 @@ void ApplicationCommonWindow::createStandardOperations()
   connect( helpAboutAction, SIGNAL( triggered() ) , this, SLOT( onAbout() ) );
   myStdActions.insert( HelpAboutId, helpAboutAction );
 
-    // popuplate a menu with all actions
+  // create preferences menu
+  QMenu* aPrefMenu = new QMenu( QObject::tr("MNU_PREFERENCES") );
+  aPrefMenu->addAction( filePrefUseVBOAction );
+  
+  // popuplate a menu with all actions
   myFilePopup = new QMenu( this );
   myFilePopup = menuBar()->addMenu( QObject::tr("MNU_FILE") );
   myFilePopup->addAction( fileNewAction );
   myFilePopup->addAction( fileCloseAction );
   myFileSeparator = myFilePopup->addSeparator();
+  myFilePopup->addMenu( aPrefMenu );
+  myFileSeparator = myFilePopup->addSeparator();
   myFilePopup->addAction( fileQuitAction );
     
        // add a view menu
@@ -124,7 +140,6 @@ void ApplicationCommonWindow::createStandardOperations()
   view->addAction( viewStatusAction );
        
        // add a help menu
-
   QMenu * help = new QMenu( this );
   menuBar()->addSeparator();
   help = menuBar()->addMenu( QObject::tr("MNU_HELP") );
@@ -136,6 +151,8 @@ void ApplicationCommonWindow::createStandardOperations()
   myStdToolBar->addAction( helpAboutAction );
        
   myStdActions.at(FileCloseId)->setEnabled(myDocuments.count() > 0);
+
+  myStdActions.at(FilePrefUseVBOId)->setEnabled( true );
 }
 
 void ApplicationCommonWindow::createCasCadeOperations()
@@ -181,9 +198,43 @@ void ApplicationCommonWindow::createCasCadeOperations()
   a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_DEL") ), QObject::tr("MNU_TOOL_DEL"), this );
   a->setToolTip( QObject::tr("TBR_TOOL_DEL") );
   a->setStatusTip( QObject::tr("TBR_TOOL_DEL") );
-  connect( a, SIGNAL( triggered() ) , this, SLOT( onToolAction() ) );
+  connect( a, SIGNAL( activated() ) , this, SLOT( onToolAction() ) );
   myToolActions.insert( ToolDeleteId, a );
 
+#ifdef HAVE_OPENCL
+
+  // populate a tool bar with some actions
+  myRaytraceBar = addToolBar( tr( "Ray-trace options" ) );
+
+  a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_SHADOWS") ), QObject::tr("MNU_TOOL_SHADOWS"), this );
+  a->setToolTip( QObject::tr("TBR_TOOL_SHADOWS") );
+  a->setStatusTip( QObject::tr("TBR_TOOL_SHADOWS") );
+  a->setCheckable( true );
+  a->setChecked( true );
+  connect( a, SIGNAL( activated() ) , this, SLOT( onRaytraceAction() ) );
+  myRaytraceActions.insert( ToolShadowsId, a );
+  myRaytraceBar->addAction( a );
+
+  a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_REFLECTIONS") ), QObject::tr("MNU_TOOL_REFLECTIONS"), this );
+  a->setToolTip( QObject::tr("TBR_TOOL_REFLECTIONS") );
+  a->setStatusTip( QObject::tr("TBR_TOOL_REFLECTIONS") );
+  a->setCheckable( true );
+  a->setChecked( true );
+  connect( a, SIGNAL( activated() ) , this, SLOT( onRaytraceAction() ) );
+  myRaytraceActions.insert( ToolReflectionsId, a );
+  myRaytraceBar->addAction( a );
+
+  a = new QAction( QPixmap( dir+QObject::tr("ICON_TOOL_ANTIALIASING") ), QObject::tr("MNU_TOOL_ANTIALIASING"), this );
+  a->setToolTip( QObject::tr("TBR_TOOL_ANTIALIASING") );
+  a->setStatusTip( QObject::tr("TBR_TOOL_ANTIALIASING") );
+  a->setCheckable( true );
+  a->setChecked( false );
+  connect( a, SIGNAL( activated() ) , this, SLOT( onRaytraceAction() ) );
+  myRaytraceActions.insert( ToolAntialiasingId, a );
+  myRaytraceBar->addAction( a );
+
+#endif
+
   QSignalMapper* sm = new QSignalMapper( this );
   connect( sm, SIGNAL( mapped( int ) ), this, SLOT( onSetMaterial( int ) ) );
 
@@ -275,12 +326,22 @@ void ApplicationCommonWindow::windowsMenuAboutToShow()
 
   QString dir = getResourceDir() + QString( "/" );
 
-  a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_NEW3D" ) ), QObject::tr( "MNU_WINDOW_NEW3D" ), this );
-  a->setToolTip( QObject::tr( "TBR_WINDOW_NEW3D" ) );
-  a->setStatusTip( QObject::tr( "TBR_WINDOW_NEW3D" ) );
+  a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_NEW3D" ) ), QObject::tr( "MNU_WINDOW_NEW3D_GL" ), this );
+  a->setToolTip( QObject::tr( "TBR_WINDOW_NEW3D_GL" ) );
+  a->setStatusTip( QObject::tr( "TBR_WINDOW_NEW3D_GL" ) );
   connect( a, SIGNAL( triggered() ), this, SLOT( onCreateNewView() ) );
   myWindowPopup->addAction( a );
 
+#ifdef HAVE_OPENCL
+
+  a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_NEW3D" ) ), QObject::tr( "MNU_WINDOW_NEW3D_RT" ), this );
+  a->setToolTip( QObject::tr( "TBR_WINDOW_NEW3D_RT" ) );
+  a->setStatusTip( QObject::tr( "TBR_WINDOW_NEW3D_RT" ) );
+  connect( a, SIGNAL( activated() ), this, SLOT( onCreateNewViewRT() ) );
+  myWindowPopup->addAction( a );
+
+#endif
+
   a = new QAction( QPixmap( dir + QObject::tr( "ICON_WINDOW_CASCADE" ) ), QObject::tr( "MNU_WINDOW_CASCADE" ), this );
   a->setToolTip( QObject::tr( "TBR_WINDOW_CASCADE" ) );
   a->setStatusTip( QObject::tr( "TBR_WINDOW_CASCADE" ) );
@@ -401,7 +462,25 @@ DocumentCommon* ApplicationCommonWindow::onNewDoc()
     return aDoc;
 }
 
-void ApplicationCommonWindow::onCloseWindow(){
+DocumentCommon* ApplicationCommonWindow::onNewDocRT()
+{
+    updateFileActions();
+    DocumentCommon* aDoc = createNewDocument();
+    aDoc->onCreateNewView(true);
+    onSelectionChanged();
+    connect( aDoc, SIGNAL( sendCloseDocument( DocumentCommon* ) ),
+             this, SLOT( onCloseDocument( DocumentCommon* ) ) );
+    connect( stWs, SIGNAL( windowActivated( QWidget* ) ),
+             this, SLOT( onWindowActivated( QWidget* ) ) );
+    connect( aDoc, SIGNAL( selectionChanged() ),
+             this, SLOT( onSelectionChanged() ) );
+    myDocuments.append( aDoc );
+    myStdActions.at( FileCloseId )->setEnabled( myDocuments.count() > 0 );
+    return aDoc;
+}
+
+void ApplicationCommonWindow::onCloseWindow()
+{
     MDIWindow* m = (MDIWindow*)stWs->activeWindow();
     if ( m )
     {
@@ -410,6 +489,27 @@ void ApplicationCommonWindow::onCloseWindow(){
     }
 }
 
+void ApplicationCommonWindow::onUseVBO()
+{
+    MDIWindow* w = ( MDIWindow* ) stWs->activeWindow();
+    
+    if ( NULL == w )
+      return;
+
+    Handle(AIS_InteractiveContext) aContextAIS = w->getDocument()->getContext();
+
+    if (aContextAIS.IsNull())
+      return;
+
+    Handle(OpenGl_GraphicDriver) aDriver =
+      Handle(OpenGl_GraphicDriver)::DownCast (aContextAIS->CurrentViewer()->Driver());
+
+    if (!aDriver.IsNull())
+    {
+      aDriver->ChangeOptions().vboDisable = Standard_True;
+    }
+}
+
 void ApplicationCommonWindow::onCloseDocument(DocumentCommon* theDoc)
 {
   myDocuments.removeAll( theDoc );
@@ -449,14 +549,38 @@ void ApplicationCommonWindow::onAbout()
 
 void ApplicationCommonWindow::onCreateNewView()
 {
-    DocumentCommon* doc = ((MDIWindow*) stWs->activeWindow())->getDocument();
-    doc->onCreateNewView();
+  MDIWindow* window = qobject_cast< MDIWindow* >( stWs->activeWindow() );
+  window->getDocument()->onCreateNewView( false );
+}
+
+#ifdef HAVE_OPENCL
+
+void ApplicationCommonWindow::onCreateNewViewRT()
+{
+  MDIWindow* window = qobject_cast< MDIWindow* >( stWs->activeWindow() );
+  window->getDocument()->onCreateNewView( true );
 }
 
+#endif
+
 void ApplicationCommonWindow::onWindowActivated ( QWidget * w )
 {
-    if ( w )
-        ((MDIWindow*) w)->onWindowActivated();
+  if (w == NULL)
+  {
+    return;
+  }
+  
+  MDIWindow* window = qobject_cast< MDIWindow* >(w);
+
+  window->onWindowActivated();
+
+#ifdef HAVE_OPENCL
+
+  myRaytraceActions.at( ToolShadowsId )->setChecked (window->ShadowsEnabled());
+  myRaytraceActions.at( ToolReflectionsId )->setChecked (window->ReflectionsEnabled());
+  myRaytraceActions.at( ToolAntialiasingId )->setChecked (window->AntialiasingEnabled());
+
+#endif
 }
 
 void ApplicationCommonWindow::onToolAction()
@@ -484,6 +608,36 @@ void ApplicationCommonWindow::onToolAction()
         doc->onDelete();
 }
 
+#ifdef HAVE_OPENCL
+
+void ApplicationCommonWindow::onRaytraceAction()
+{
+  QAction* sentBy = (QAction*) sender();
+  
+  DocumentCommon* doc = qobject_cast< MDIWindow* >(
+    ApplicationCommonWindow::getWorkspace()->activeWindow())->getDocument();
+
+  if( sentBy == myRaytraceActions.at( ToolShadowsId ) )
+  {
+    bool flag = myRaytraceActions.at( ToolShadowsId )->isChecked(); 
+    doc->onShadows( flag );
+  }
+
+  if( sentBy == myRaytraceActions.at( ToolReflectionsId ) )
+  {
+    bool flag = myRaytraceActions.at( ToolReflectionsId )->isChecked();
+    doc->onReflections( flag );
+  }
+
+  if( sentBy == myRaytraceActions.at( ToolAntialiasingId ) )
+  {
+    bool flag = myRaytraceActions.at( ToolAntialiasingId )->isChecked();
+    doc->onAntialiasing( flag );
+  }
+}
+
+#endif
+
 void ApplicationCommonWindow::onSelectionChanged()
 {
   QWorkspace* ws = ApplicationCommonWindow::getWorkspace();
index 742af43..787f458 100755 (executable)
 #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                            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;
index 806deda..a5be1b3 100755 (executable)
         <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>
     </message>
         <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>
index 86fff36..cd87a19 100755 (executable)
         <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>
     </message>
         <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>
     </message>
         <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>
     </message>
     </message>
     <message>
         <source>TBR_WINDOW_NEW3D</source>
-        <translation>New 3d View</translation>
+        <translation>New 3D View</translation>
+    </message>
+    <message>
+        <source>TBR_WINDOW_NEW3D_GL</source>
+        <translation>New GL 3D View</translation>
+    </message>
+    <message>
+        <source>MNU_WINDOW_NEW3D_GL</source>
+        <translation>New GL 3D View</translation>
+    </message>
+    <message>
+        <source>TBR_WINDOW_NEW3D_RT</source>
+        <translation>&amp;New RT 3D View</translation>
     </message>
     <message>
-        <source>MNU_WINDOW_NEW3D</source>
-        <translation>&amp;New 3d View</translation>
+        <source>MNU_WINDOW_NEW3D_RT</source>
+        <translation>&amp;New RT 3D View</translation>
     </message>
     <message>
         <source>MNU_STATUS_BAR</source>
index d8de9df..71647c8 100755 (executable)
@@ -50,12 +50,10 @@ myNbViews( 0 )
   myViewer = Viewer( getenv("DISPLAY"), a3DName.ToExtString(), "", 1000.0,
                               V3d_XposYnegZpos, Standard_True, Standard_True );
 
-  myViewer->Init();
        myViewer->SetDefaultLights();
        myViewer->SetLightOn();
 
-       myContext =new AIS_InteractiveContext(myViewer);        
-       //onCreateNewView();
+       myContext =new AIS_InteractiveContext(myViewer);
 }
 
 DocumentCommon::~DocumentCommon()
@@ -67,15 +65,16 @@ ApplicationCommonWindow* DocumentCommon::getApplication()
        return myApp;
 }
 
-MDIWindow* DocumentCommon::createNewMDIWindow()
+MDIWindow* DocumentCommon::createNewMDIWindow( bool theRT )
 {
   QWorkspace* ws = myApp->getWorkspace();
-  return new MDIWindow( this, ws, 0 );
+  return new MDIWindow( this, ws, 0, theRT );
 }
-void DocumentCommon::onCreateNewView() 
+
+void DocumentCommon::onCreateNewView( bool theRT ) 
 {
        QWorkspace* ws = myApp->getWorkspace();
-  MDIWindow* w = createNewMDIWindow();
+  MDIWindow* w = createNewMDIWindow( theRT );
        
        if( !w )
          return;
@@ -228,4 +227,48 @@ void DocumentCommon::onDelete()
     getApplication()->onSelectionChanged();
 }
 
+#ifdef HAVE_OPENCL
+
+void DocumentCommon::onShadows( int state )
+{
+  QWorkspace* ws = ApplicationCommonWindow::getWorkspace();
+
+  MDIWindow* window = qobject_cast< MDIWindow* >( ws->activeWindow() );
+
+  if( window == NULL )
+    return;
+
+  window->setRaytracedShadows( state );
+
+  myContext->UpdateCurrentViewer();
+}
+
+void DocumentCommon::onReflections( int state )
+{
+  QWorkspace* ws = ApplicationCommonWindow::getWorkspace();
+
+  MDIWindow* window = qobject_cast< MDIWindow* >( ws->activeWindow() );
+
+  if( window == NULL )
+    return;
 
+  window->setRaytracedReflections( state );
+
+  myContext->UpdateCurrentViewer();
+}
+
+void DocumentCommon::onAntialiasing( int state )
+{
+  QWorkspace* ws = ApplicationCommonWindow::getWorkspace();
+
+  MDIWindow* window = qobject_cast< MDIWindow* >( ws->activeWindow() );
+
+  if( window == NULL )
+    return;
+
+  window->setRaytracedAntialiasing( state );
+
+  myContext->UpdateCurrentViewer();
+}
+
+#endif
index 3dcefb4..b8347fc 100755 (executable)
@@ -28,7 +28,7 @@ public:
        void                           fitAll();
        
 protected:
-  virtual MDIWindow*                   createNewMDIWindow();
+  virtual MDIWindow*                   createNewMDIWindow( bool theRT = false );
 
 signals:
   void                           selectionChanged();
@@ -36,10 +36,17 @@ signals:
 
 public slots:
        virtual void                   onCloseView( MDIWindow* );
-       virtual void                   onCreateNewView();
+       virtual void                   onCreateNewView( bool theRT = false );
        virtual void                   onMaterial();
-        virtual void                   onMaterial( int );
+  virtual void                   onMaterial( int );
        virtual void                   onDelete();
+
+#ifdef HAVE_OPENCL
+  virtual void                   onShadows( int state );
+  virtual void                   onReflections( int state );
+  virtual void                   onAntialiasing( int state );
+#endif
+
        void                           onWireframe();
        void                           onShading();
        void                           onColor();
index 86bc793..d44da21 100755 (executable)
@@ -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
+
 
index 59e7df8..e2753cc 100755 (executable)
@@ -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
index a881a83..920face 100755 (executable)
 #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 );
-               QAction* a = new QAction( QObject::tr("MNU_CH_BACK"), this );
-               a->setToolTip( QObject::tr("TBR_CH_BACK") );
-    connect( a,SIGNAL( triggered() ), this, SLOT( onBackground() ) );
-               myBackMenu->addAction( a );  
-    addItemInPopup(myBackMenu);
+    if (!myBackMenu)
+    {
+      myBackMenu = new QMenu( 0 );
+
+                 QAction* a = new QAction( QObject::tr("MNU_CH_BACK"), this );
+                 a->setToolTip( QObject::tr("TBR_CH_BACK") );
+      connect( a, SIGNAL( activated() ), this, SLOT( onBackground() ) );
+                 myBackMenu->addAction( a );  
+      addItemInPopup(myBackMenu);
+
+      a = new QAction( QObject::tr("MNU_CH_ENV_MAP"), this );
+                 a->setToolTip( QObject::tr("TBR_CH_ENV_MAP") );
+      connect( a, SIGNAL( activated() ), this, SLOT( onEnvironmentMap() ) );
+      a->setCheckable( true );
+      a->setChecked( false );
+                 myBackMenu->addAction( a );  
+      addItemInPopup(myBackMenu);
+    }
+
     myBackMenu->exec( QCursor::pos() );
-    delete myBackMenu;
   }
   if ( w )
     w->setFocus();
@@ -946,6 +993,26 @@ void View::onBackground()
     myView->Redraw();
 }
 
+void View::onEnvironmentMap()
+{
+  if (myBackMenu->actions().at(1)->isChecked())
+  {
+    QString fileName = QFileDialog::getOpenFileName(this, tr("Open File"), "",
+                           tr("All Image Files (*.bmp *.gif *.jpg *.jpeg *.png *.tga)"));
+
+    Handle(Graphic3d_TextureEnv) aTexture = new Graphic3d_TextureEnv( fileName.toAscii().data() );
+
+    myView->SetTextureEnv (aTexture);
+    myView->SetSurfaceDetail (V3d_TEX_ENVIRONMENT);
+  }
+  else
+  {
+    myView->SetSurfaceDetail (V3d_TEX_NONE);
+  }
+  
+  myView->Redraw();
+}
+
 bool View::dump(Standard_CString theFile)
 {
   myView->Redraw();
index 50c4a7c..dac41c3 100755 (executable)
@@ -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
 };
 
index 70a1667..0cad367 100644 (file)
@@ -16,4 +16,4 @@ Global
        GlobalSection(SolutionProperties) = preSolution
                HideSolutionNode = FALSE
        EndGlobalSection
-EndGlobal
+EndGlobal
\ No newline at end of file
index 93b524d..f2b46ec 100755 (executable)
@@ -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);
index 21580fc..a586f17 100755 (executable)
@@ -6,6 +6,7 @@ CSF_objc
 CSF_Appkit
 CSF_IOKit
 CSF_OpenGlLibs
+CSF_OPENCL
 CSF_AviLibs
 CSF_FREETYPE
 CSF_GL2PS
index 57ae623..a1686f8 100755 (executable)
@@ -133,3 +133,11 @@ OpenGl_ShaderStates.cxx
 Handle_OpenGl_ShaderObject.hxx
 Handle_OpenGl_ShaderProgram.hxx
 Handle_OpenGl_ShaderManager.hxx
+OpenGl_Cl.hxx
+OpenGl_AABB.hxx
+OpenGl_AABB.cxx
+OpenGl_SceneGeometry.hxx
+OpenGl_SceneGeometry.cxx
+OpenGl_RaytraceTypes.hxx
+OpenGl_RaytraceSource.cxx
+OpenGl_Workspace_Raytrace.cxx
\ No newline at end of file
diff --git a/src/OpenGl/OpenGl_AABB.cxx b/src/OpenGl/OpenGl_AABB.cxx
new file mode 100644 (file)
index 0000000..6dcdc2b
--- /dev/null
@@ -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
\ No newline at end of file
diff --git a/src/OpenGl/OpenGl_AABB.hxx b/src/OpenGl/OpenGl_AABB.hxx
new file mode 100644 (file)
index 0000000..9893406
--- /dev/null
@@ -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
index 39c14a8..5b57f7b 100644 (file)
@@ -35,7 +35,8 @@ OpenGl_Caps::OpenGl_Caps()
 #else
   contextDebug      (Standard_False),
 #endif
-  contextNoAccel    (Standard_False)
+  contextNoAccel    (Standard_False),
+  keepArrayData     (Standard_False)
 {
   //
 }
index 4ca46fc..a7935bf 100644 (file)
@@ -67,6 +67,13 @@ public: //! @name context creation parameters
    */
   Standard_Boolean contextNoAccel;
 
+  /**
+   * Disables freeing CPU memory after building VBOs.
+   *
+   * OFF by default.
+   */
+  Standard_Boolean keepArrayData;
+
 public: //! @name class methods
 
   //! Default constructor - initialize with most optimal values.
diff --git a/src/OpenGl/OpenGl_Cl.hxx b/src/OpenGl/OpenGl_Cl.hxx
new file mode 100644 (file)
index 0000000..2bfcf33
--- /dev/null
@@ -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__
index 5b04b88..5465b52 100644 (file)
@@ -23,7 +23,7 @@
 #include <OpenGl_Display.hxx>
 
 #ifdef HAVE_CONFIG_H
-# include <config.h>
+  #include <config.h>
 #endif
 
 #ifdef HAVE_GL2PS
index be8fcd7..18ca913 100755 (executable)
 // 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>
@@ -155,6 +160,34 @@ Standard_Boolean OpenGl_GraphicDriver::SetImmediateModeDrawToFront (const Graphi
 }
 
 // =======================================================================
+// 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  :
 // =======================================================================
index 39415f6..a09ad54 100644 (file)
@@ -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.
index c1d9f57..68d3aaa 100755 (executable)
@@ -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)
   {
index 5c6a63c..b1151fb 100644 (file)
 // purpose or non-infringement. Please see the License for the specific terms
 // and conditions governing the rights and limitations under the License.
 
-#include <OpenGl_Group.hxx>
 
+#ifdef HAVE_CONFIG_H
+  #include <config.h>
+#endif
+
+#include <OpenGl_Group.hxx>
 #include <OpenGl_PrimitiveArray.hxx>
+#include <OpenGl_Structure.hxx>
 #include <OpenGl_Workspace.hxx>
 
 // =======================================================================
 // function : OpenGl_Group
 // purpose  :
 // =======================================================================
-OpenGl_Group::OpenGl_Group ()
-: myAspectLine (NULL),
-  myAspectFace (NULL),
-  myAspectMarker (NULL),
-  myAspectText (NULL),
-  myFirst (NULL),
-  myLast (NULL)
+#ifndef HAVE_OPENCL
+OpenGl_Group::OpenGl_Group()
+#else
+OpenGl_Group::OpenGl_Group (const OpenGl_Structure* theAncestorStructure)
+#endif
+: myAspectLine(NULL),
+  myAspectFace(NULL),
+  myAspectMarker(NULL),
+  myAspectText(NULL),
+  myFirst(NULL),
+  myLast(NULL)
 {
+#ifdef HAVE_OPENCL
+  myAncestorStructure = theAncestorStructure;
+  myIsRaytracable = Standard_False;
+  myModificationState = 0; // initial state
+#endif
 }
 
 // =======================================================================
@@ -89,6 +103,18 @@ void OpenGl_Group::SetAspectFace (const CALL_DEF_CONTEXTFILLAREA& theAspect,
     anAspectFace->SetAspect (theAspect);
     AddElement (TelNil/*TelAspectFace*/, anAspectFace);
   }
+
+#ifdef HAVE_OPENCL
+  if (myIsRaytracable)
+  {
+    myModificationState++;
+
+    if (myAncestorStructure != NULL)
+    {
+      myAncestorStructure->UpdateStateWithAncestorStructures();
+    }
+  }
+#endif
 }
 
 // =======================================================================
@@ -141,15 +167,29 @@ void OpenGl_Group::SetAspectText (const CALL_DEF_CONTEXTTEXT& theAspect,
 // function : AddElement
 // purpose  :
 // =======================================================================
-void OpenGl_Group::AddElement (const TelType AType, OpenGl_Element *AElem )
+void OpenGl_Group::AddElement (const TelType theType, OpenGl_Element *theElem)
 {
-  OpenGl_ElementNode *node = new OpenGl_ElementNode();
+  OpenGl_ElementNode *aNode = new OpenGl_ElementNode();
+
+  aNode->type = theType;
+  aNode->elem = theElem;
+  aNode->next = NULL;
+  (myLast? myLast->next : myFirst) = aNode;
+  myLast = aNode;
+
+#ifdef HAVE_OPENCL
+  if (OpenGl_Raytrace::IsRaytracedElement (aNode))
+  {
+    myModificationState++;
+    myIsRaytracable = Standard_True;
 
-  node->type = AType;
-  node->elem = AElem;
-  node->next = NULL;
-  (myLast? myLast->next : myFirst) = node;
-  myLast = node;
+    if (myAncestorStructure != NULL)
+    {
+      myAncestorStructure->UpdateStateWithAncestorStructures();
+      myAncestorStructure->SetRaytracableWithAncestorStructures();
+    }
+  }
+#endif
 }
 
 // =======================================================================
index 6ed3e99..85c085c 100644 (file)
@@ -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,19 +65,41 @@ public:
   virtual void Render  (const Handle(OpenGl_Workspace)& theWorkspace) const;
   virtual void Release (const Handle(OpenGl_Context)&   theGlCtx);
 
+  //! Returns first OpenGL element node of the group.
+  const OpenGl_ElementNode* FirstNode() const { return myFirst; }
+
+  //! Returns OpenGL face aspect.
+  const OpenGl_AspectFace* AspectFace() const { return myAspectFace; }
+
+#ifdef HAVE_OPENCL
+
+  //! Returns modification state for ray-tracing.
+  Standard_Size ModificationState() const { return myModificationState; }
+
+  //! Is the group ray-tracable (contains ray-tracable elements)?
+  Standard_Boolean IsRaytracable() const { return myIsRaytracable; }
+
+#endif
+
 protected:
 
   virtual ~OpenGl_Group();
 
 protected:
 
-  OpenGl_AspectLine*   myAspectLine;
-  OpenGl_AspectFace*   myAspectFace;
-  OpenGl_AspectMarker* myAspectMarker;
-  OpenGl_AspectText*   myAspectText;
-
-  OpenGl_ElementNode*  myFirst;
-  OpenGl_ElementNode*  myLast;
+  OpenGl_AspectLine*     myAspectLine;
+  OpenGl_AspectFace*     myAspectFace;
+  OpenGl_AspectMarker*   myAspectMarker;
+  OpenGl_AspectText*     myAspectText;
+
+  OpenGl_ElementNode*    myFirst;
+  OpenGl_ElementNode*    myLast;
+  
+#ifdef HAVE_OPENCL
+  const OpenGl_Structure*  myAncestorStructure;
+  Standard_Boolean         myIsRaytracable;
+  Standard_Size            myModificationState;
+#endif
 
 public:
 
index cc25f07..2e4b992 100644 (file)
 // 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;
     }
   }
index eceae34..73828ec 100644 (file)
@@ -80,7 +80,17 @@ class OpenGl_LayerList
   
   //! Render this element
   void Render (const Handle(OpenGl_Workspace) &theWorkspace) const;
+
+  //! Returns the set of OpenGL Z-layers.
+  const OpenGl_SequenceOfLayers& Layers() const { return myLayers; }
   
+#ifdef HAVE_OPENCL
+
+  //! Returns structure modification state (for ray-tracing).
+  Standard_Size ModificationState() const { return myModificationState; }
+
+#endif
+
  private:
   
   //! Get default layer
@@ -94,6 +104,10 @@ class OpenGl_LayerList
   Standard_Integer        myNbPriorities;
   Standard_Integer        myNbStructures;
 
+#ifdef HAVE_OPENCL
+  mutable Standard_Size   myModificationState;
+#endif
+
  public:
   DEFINE_STANDARD_ALLOC
 };
index ceda9b7..0a5b0b2 100755 (executable)
@@ -149,7 +149,11 @@ Standard_Boolean OpenGl_PrimitiveArray::BuildVBO (const Handle(OpenGl_Workspace)
     }
   }
 
-  clearMemoryOwn();
+  if (!aGlCtx->caps->keepArrayData)
+  {
+    clearMemoryOwn();
+  }
+  
   return Standard_True;
 }
 
index 8ea8160..7e43b7c 100644 (file)
@@ -58,6 +58,9 @@ class OpenGl_PriorityList
   //! or less). Returns Standard_False if the list can not be accepted.
   Standard_Boolean Append (const OpenGl_PriorityList& theOther);
 
+  //! Returns array of OpenGL structures.
+  const OpenGl_ArrayOfStructure& ArrayOfStructures() const { return myArray; }
+
  protected:
 
   OpenGl_ArrayOfStructure myArray;
diff --git a/src/OpenGl/OpenGl_RaytraceSource.cxx b/src/OpenGl/OpenGl_RaytraceSource.cxx
new file mode 100644 (file)
index 0000000..baf9104
--- /dev/null
@@ -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
diff --git a/src/OpenGl/OpenGl_RaytraceTypes.hxx b/src/OpenGl/OpenGl_RaytraceTypes.hxx
new file mode 100644 (file)
index 0000000..c75ae3f
--- /dev/null
@@ -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
diff --git a/src/OpenGl/OpenGl_SceneGeometry.cxx b/src/OpenGl/OpenGl_SceneGeometry.cxx
new file mode 100644 (file)
index 0000000..5184e8d
--- /dev/null
@@ -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
diff --git a/src/OpenGl/OpenGl_SceneGeometry.hxx b/src/OpenGl/OpenGl_SceneGeometry.hxx
new file mode 100644 (file)
index 0000000..d01af3c
--- /dev/null
@@ -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
index 5f9c9d7..42ade61 100644 (file)
 // 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;
@@ -307,28 +335,150 @@ void OpenGl_Structure::ClearHighlightColor (const Handle(OpenGl_Context)& theGlC
 }
 
 // =======================================================================
+// function : SetNamedStatus
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::SetNamedStatus (const Standard_Integer aStatus)
+{
+  myNamedStatus = aStatus;
+
+#ifdef HAVE_OPENCL
+  if (myIsRaytracable)
+  {
+    UpdateStateWithAncestorStructures();
+  }
+#endif
+}
+
+#ifdef HAVE_OPENCL
+
+// =======================================================================
+// function : RegisterAncestorStructure
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::RegisterAncestorStructure (const OpenGl_Structure* theStructure) const
+{
+  for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+  {
+    if (anIt.Value() == theStructure)
+    {
+      return;
+    }    
+  }
+
+  myAncestorStructures.Append (theStructure);
+}
+
+// =======================================================================
+// function : UnregisterAncestorStructure
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::UnregisterAncestorStructure (const OpenGl_Structure* theStructure) const
+{
+  for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+  {
+    if (anIt.Value() == theStructure)
+    {
+      myAncestorStructures.Remove (anIt);
+      return;
+    }    
+  }
+}
+
+// =======================================================================
+// function : UpdateStateWithAncestorStructures
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::UpdateStateWithAncestorStructures() const
+{
+  myModificationState++;
+
+  for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+  {
+    anIt.Value()->UpdateStateWithAncestorStructures();
+  }
+}
+
+// =======================================================================
+// function : UpdateRaytracableWithAncestorStructures
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::UpdateRaytracableWithAncestorStructures() const
+{
+  myIsRaytracable = OpenGl_Raytrace::IsRaytracedStructure (this);
+
+  if (!myIsRaytracable)
+  {
+    for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+    {
+      anIt.Value()->UpdateRaytracableWithAncestorStructures();
+    }
+  }
+}
+
+// =======================================================================
+// function : SetRaytracableWithAncestorStructures
+// purpose  :
+// =======================================================================
+void OpenGl_Structure::SetRaytracableWithAncestorStructures() const
+{
+  myIsRaytracable = Standard_True;
+
+  for (OpenGl_ListOfStructure::Iterator anIt (myAncestorStructures); anIt.More(); anIt.Next())
+  {
+    if (!anIt.Value()->IsRaytracable())
+    {
+      anIt.Value()->SetRaytracableWithAncestorStructures();
+    }
+  }
+}
+
+#endif
+
+// =======================================================================
 // function : Connect
 // purpose  :
 // =======================================================================
-void OpenGl_Structure::Connect (const OpenGl_Structure *AStructure)
+void OpenGl_Structure::Connect (const OpenGl_Structure *theStructure)
 {
-  Disconnect (AStructure);
-  myConnected.Append(AStructure);
+  Disconnect (theStructure);
+  myConnected.Append (theStructure);
+
+#ifdef HAVE_OPENCL
+  if (theStructure->IsRaytracable())
+  {
+    UpdateStateWithAncestorStructures();
+    SetRaytracableWithAncestorStructures();
+  }
+
+  theStructure->RegisterAncestorStructure (this);
+#endif
 }
 
 // =======================================================================
 // function : Disconnect
 // purpose  :
 // =======================================================================
-void OpenGl_Structure::Disconnect (const OpenGl_Structure *AStructure)
+void OpenGl_Structure::Disconnect (const OpenGl_Structure *theStructure)
 {
-  OpenGl_ListOfStructure::Iterator its(myConnected);
+  OpenGl_ListOfStructure::Iterator its (myConnected);
   while (its.More())
   {
     // Check for the given structure
-    if (its.Value() == AStructure)
+    if (its.Value() == theStructure)
     {
-      myConnected.Remove(its);
+      myConnected.Remove (its);
+
+#ifdef HAVE_OPENCL
+      if (theStructure->IsRaytracable())
+      {
+        UpdateStateWithAncestorStructures();
+        UpdateRaytracableWithAncestorStructures();
+      }
+
+      theStructure->UnregisterAncestorStructure (this);
+#endif
+
       return;
     }
     its.Next();
@@ -339,10 +489,15 @@ void OpenGl_Structure::Disconnect (const OpenGl_Structure *AStructure)
 // function : AddGroup
 // purpose  :
 // =======================================================================
-OpenGl_Group * OpenGl_Structure::AddGroup ()
+OpenGl_Group * OpenGl_Structure::AddGroup()
 {
   // Create new group
-  OpenGl_Group *g = new OpenGl_Group;
+#ifndef HAVE_OPENCL
+  OpenGl_Group *g = new OpenGl_Group();
+#else
+  OpenGl_Group *g = new OpenGl_Group (this);
+#endif
+
   myGroups.Append(g);
   return g;
 }
@@ -359,9 +514,18 @@ void OpenGl_Structure::RemoveGroup (const Handle(OpenGl_Context)& theGlCtx,
     // Check for the given group
     if (anIter.Value() == theGroup)
     {
-      // Delete object
-      OpenGl_Element::Destroy (theGlCtx, const_cast<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
 }
 
 // =======================================================================
index 8d49896..994e6ad 100644 (file)
@@ -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
index 32e4b2d..29244ef 100644 (file)
 // 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
 }
 
 /*----------------------------------------------------------------------*/
@@ -165,8 +174,21 @@ void OpenGl_View::SetTextureEnv (const Handle(OpenGl_Context)&       theCtx,
 
   myTextureEnv = new OpenGl_Texture (theTexture->GetParams());
   Handle(Image_PixMap) anImage = theTexture->GetImage();
-  if( !anImage.IsNull())
+  if (!anImage.IsNull())
     myTextureEnv->Init (theCtx, *anImage.operator->(), theTexture->Type());
+
+#ifdef HAVE_OPENCL
+  myModificationState++;
+#endif
+}
+
+void OpenGl_View::SetSurfaceDetail (const Visual3d_TypeOfSurfaceDetail theMode)
+{
+  mySurfaceDetail = theMode;
+
+#ifdef HAVE_OPENCL
+  myModificationState++;
+#endif
 }
 
 /*----------------------------------------------------------------------*/
index fb1636d..6419c1c 100644 (file)
@@ -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
 };
index 7200c3a..c11b1f9 100644 (file)
@@ -636,73 +636,14 @@ call_util_mat_mul( matrix3 mat_a, matrix3 mat_b, matrix3 mat_c)
 
 /*----------------------------------------------------------------------*/
 
-//call_func_redraw_all_structs_proc
-void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext,
-                          const Handle(OpenGl_Workspace) &AWorkspace,
-                          const Graphic3d_CView& ACView,
-                          const Aspect_CLayer2d& ACUnderLayer,
-                          const Aspect_CLayer2d& ACOverLayer)
+void OpenGl_View::DrawBackground (const Handle(OpenGl_Workspace) &AWorkspace)
 {
-  // Store and disable current clipping planes
-  const Handle(OpenGl_Context)& aContext = AWorkspace->GetGlContext();
-  const Standard_Integer aMaxClipPlanes = aContext->MaxClipPlanes();
-  const GLenum lastid = GL_CLIP_PLANE0 + aMaxClipPlanes;
-  OPENGL_CLIP_PLANE *oldPlanes = new OPENGL_CLIP_PLANE[aMaxClipPlanes];
-  OPENGL_CLIP_PLANE *ptrPlane = oldPlanes;
-  GLenum planeid = GL_CLIP_PLANE0;
-  for ( ; planeid < lastid; planeid++, ptrPlane++ )
-  {
-    glGetClipPlane( planeid, ptrPlane->Equation );
-    if ( ptrPlane->isEnabled )
-    {
-      glDisable( planeid );
-      ptrPlane->isEnabled = GL_TRUE;
-    }
-    else
-    {
-      ptrPlane->isEnabled = GL_FALSE;
-    }
-  }
-
-  // Set OCCT state uniform variables
-
-  if (!aContext->ShaderManager()->IsEmpty())
-  {
-    if (myLightSourcesChanged)
-    {
-      aContext->ShaderManager()->UpdateLightSourceStateTo (&myLights);
-      myLightSourcesChanged = Standard_False;
-    }
-
-    if (myViewMappingChanged)
-    {
-      aContext->ShaderManager()->UpdateProjectionStateTo (myMappingMatrix);
-      myViewMappingChanged = Standard_False;
-    }
-
-    if (myOrientationChanged)
-    {
-      aContext->ShaderManager()->UpdateWorldViewStateTo (myOrientationMatrix);
-      myOrientationChanged = Standard_False;
-    }
-
-    if (aContext->ShaderManager()->ModelWorldState().Index() == 0)
-    {
-      Tmatrix3 aModelWorldState = { { 1.f, 0.f, 0.f, 0.f },
-                                    { 0.f, 1.f, 0.f, 0.f },
-                                    { 0.f, 0.f, 1.f, 0.f },
-                                    { 0.f, 0.f, 0.f, 1.f } };
-
-      aContext->ShaderManager()->UpdateModelWorldStateTo (aModelWorldState);
-    }
-  }
-
   /////////////////////////////////////////////////////////////////////////////
   // Step 1: Prepare for redraw
 
   // Render background
   if ( (AWorkspace->NamedStatus & OPENGL_NS_WHITEBACK) == 0 &&
-          ( myBgTexture.TexId != 0 || myBgGradient.type != Aspect_GFM_NONE ) )
+    ( myBgTexture.TexId != 0 || myBgGradient.type != Aspect_GFM_NONE ) )
   {
     const Standard_Integer aViewWidth = AWorkspace->Width();
     const Standard_Integer aViewHeight = AWorkspace->Height();
@@ -723,8 +664,8 @@ void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext,
     // - gradient fill type is not Aspect_GFM_NONE and
     // - either background texture is no specified or it is drawn in Aspect_FM_CENTERED mode
     if ( ( myBgGradient.type != Aspect_GFM_NONE ) &&
-         ( myBgTexture.TexId == 0 || myBgTexture.Style == Aspect_FM_CENTERED ||
-           myBgTexture.Style == Aspect_FM_NONE ) )
+      ( myBgTexture.TexId == 0 || myBgTexture.Style == Aspect_FM_CENTERED ||
+      myBgTexture.Style == Aspect_FM_NONE ) )
     {
       Tfloat* corner1 = 0;/* -1,-1*/
       Tfloat* corner2 = 0;/*  1,-1*/
@@ -735,63 +676,63 @@ void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext,
 
       switch( myBgGradient.type )
       {
-        case Aspect_GFM_HOR:
-          corner1 = myBgGradient.color1.rgb;
-          corner2 = myBgGradient.color2.rgb;
-          corner3 = myBgGradient.color2.rgb;
-          corner4 = myBgGradient.color1.rgb;
-          break;
-        case Aspect_GFM_VER:
-          corner1 = myBgGradient.color2.rgb;
-          corner2 = myBgGradient.color2.rgb;
-          corner3 = myBgGradient.color1.rgb;
-          corner4 = myBgGradient.color1.rgb;
-          break;
-        case Aspect_GFM_DIAG1:
-          corner2 = myBgGradient.color2.rgb;
-          corner4 = myBgGradient.color1.rgb;
-          dcorner1 [0] = dcorner2[0] = 0.5F * (corner2[0] + corner4[0]);
-          dcorner1 [1] = dcorner2[1] = 0.5F * (corner2[1] + corner4[1]);
-          dcorner1 [2] = dcorner2[2] = 0.5F * (corner2[2] + corner4[2]);
-          corner1 = dcorner1;
-          corner3 = dcorner2;
-          break;
-        case Aspect_GFM_DIAG2:
-          corner1 = myBgGradient.color2.rgb;
-          corner3 = myBgGradient.color1.rgb;
-          dcorner1 [0] = dcorner2[0] = 0.5F * (corner1[0] + corner3[0]);
-          dcorner1 [1] = dcorner2[1] = 0.5F * (corner1[1] + corner3[1]);
-          dcorner1 [2] = dcorner2[2] = 0.5F * (corner1[2] + corner3[2]);
-          corner2 = dcorner1;
-          corner4 = dcorner2;
-          break;
-        case Aspect_GFM_CORNER1:
-          corner1 = myBgGradient.color2.rgb;
-          corner2 = myBgGradient.color2.rgb;
-          corner3 = myBgGradient.color2.rgb;
-          corner4 = myBgGradient.color1.rgb;
-          break;
-        case Aspect_GFM_CORNER2:
-          corner1 = myBgGradient.color2.rgb;
-          corner2 = myBgGradient.color2.rgb;
-          corner3 = myBgGradient.color1.rgb;
-          corner4 = myBgGradient.color2.rgb;
-          break;
-        case Aspect_GFM_CORNER3:
-          corner1 = myBgGradient.color2.rgb;
-          corner2 = myBgGradient.color1.rgb;
-          corner3 = myBgGradient.color2.rgb;
-          corner4 = myBgGradient.color2.rgb;
-          break;
-        case Aspect_GFM_CORNER4:
-          corner1 = myBgGradient.color1.rgb;
-          corner2 = myBgGradient.color2.rgb;
-          corner3 = myBgGradient.color2.rgb;
-          corner4 = myBgGradient.color2.rgb;
-          break;
-        default:
-          //printf("gradient background type not right\n");
-         break;
+      case Aspect_GFM_HOR:
+        corner1 = myBgGradient.color1.rgb;
+        corner2 = myBgGradient.color2.rgb;
+        corner3 = myBgGradient.color2.rgb;
+        corner4 = myBgGradient.color1.rgb;
+        break;
+      case Aspect_GFM_VER:
+        corner1 = myBgGradient.color2.rgb;
+        corner2 = myBgGradient.color2.rgb;
+        corner3 = myBgGradient.color1.rgb;
+        corner4 = myBgGradient.color1.rgb;
+        break;
+      case Aspect_GFM_DIAG1:
+        corner2 = myBgGradient.color2.rgb;
+        corner4 = myBgGradient.color1.rgb;
+        dcorner1 [0] = dcorner2[0] = 0.5F * (corner2[0] + corner4[0]);
+        dcorner1 [1] = dcorner2[1] = 0.5F * (corner2[1] + corner4[1]);
+        dcorner1 [2] = dcorner2[2] = 0.5F * (corner2[2] + corner4[2]);
+        corner1 = dcorner1;
+        corner3 = dcorner2;
+        break;
+      case Aspect_GFM_DIAG2:
+        corner1 = myBgGradient.color2.rgb;
+        corner3 = myBgGradient.color1.rgb;
+        dcorner1 [0] = dcorner2[0] = 0.5F * (corner1[0] + corner3[0]);
+        dcorner1 [1] = dcorner2[1] = 0.5F * (corner1[1] + corner3[1]);
+        dcorner1 [2] = dcorner2[2] = 0.5F * (corner1[2] + corner3[2]);
+        corner2 = dcorner1;
+        corner4 = dcorner2;
+        break;
+      case Aspect_GFM_CORNER1:
+        corner1 = myBgGradient.color2.rgb;
+        corner2 = myBgGradient.color2.rgb;
+        corner3 = myBgGradient.color2.rgb;
+        corner4 = myBgGradient.color1.rgb;
+        break;
+      case Aspect_GFM_CORNER2:
+        corner1 = myBgGradient.color2.rgb;
+        corner2 = myBgGradient.color2.rgb;
+        corner3 = myBgGradient.color1.rgb;
+        corner4 = myBgGradient.color2.rgb;
+        break;
+      case Aspect_GFM_CORNER3:
+        corner1 = myBgGradient.color2.rgb;
+        corner2 = myBgGradient.color1.rgb;
+        corner3 = myBgGradient.color2.rgb;
+        corner4 = myBgGradient.color2.rgb;
+        break;
+      case Aspect_GFM_CORNER4:
+        corner1 = myBgGradient.color1.rgb;
+        corner2 = myBgGradient.color2.rgb;
+        corner3 = myBgGradient.color2.rgb;
+        corner4 = myBgGradient.color2.rgb;
+        break;
+      default:
+        //printf("gradient background type not right\n");
+        break;
       }
 
       // Save GL parameters
@@ -831,7 +772,7 @@ void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext,
       GLfloat texX_range = 1.F; // texture <s> coordinate
       GLfloat texY_range = 1.F; // texture <t> coordinate
 
-         // Set up for stretching or tiling
+      // Set up for stretching or tiling
       GLfloat x_offset, y_offset;
       if ( myBgTexture.Style == Aspect_FM_CENTERED )
       {
@@ -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;
index e7562f9..e6fdf67 100644 (file)
 // 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,11 +507,29 @@ void OpenGl_Workspace::Redraw (const Graphic3d_CView& theCView,
     toSwap = 0; // no need to swap buffers
   }
 
-  Redraw1 (theCView, theCUnderLayer, theCOverLayer, toSwap);
-  if (aFrameBuffer == NULL || !myTransientDrawToFront)
+#ifdef HAVE_OPENCL
+  if (!theCView.IsRaytracing || myComputeInitStatus == OpenGl_CLIS_FAIL)
+  {
+#endif
+    Redraw1 (theCView, theCUnderLayer, theCOverLayer, toSwap);
+    if (aFrameBuffer == NULL || !myTransientDrawToFront)
+    {
+      RedrawImmediatMode();
+    }
+
+    theCView.WasRedrawnGL = Standard_True;
+#ifdef HAVE_OPENCL
+  }
+  else
   {
-    RedrawImmediatMode();
+    int aSizeX = aFrameBuffer != NULL ? aFrameBuffer->GetVPSizeX() : myWidth;
+    int aSizeY = aFrameBuffer != NULL ? aFrameBuffer->GetVPSizeY() : myHeight;
+
+    Raytrace (theCView, aSizeX, aSizeY, toSwap);
+
+    theCView.WasRedrawnGL = Standard_False;
   }
+#endif
 
   if (aFrameBuffer != NULL)
   {
index c17a006..4f13dbd 100644 (file)
 #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.
-//! Provides methods to render primitives and maintan GL state.
+//! Represents window with GL context.
+//! Provides methods to render primitives and maintain GL state.
 class OpenGl_Workspace : public OpenGl_Window
 {
 public:
@@ -208,6 +218,171 @@ protected:
   void setTextureParams (Handle(OpenGl_Texture)&                theTexture,
                          const Handle(Graphic3d_TextureParams)& theParams);
 
+#ifdef HAVE_OPENCL
+
+public:
+
+  //! Returns information about OpenCL device used for computations.
+  Standard_Boolean GetOpenClDeviceInfo (
+    NCollection_DataMap<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;
diff --git a/src/OpenGl/OpenGl_Workspace_Raytrace.cxx b/src/OpenGl/OpenGl_Workspace_Raytrace.cxx
new file mode 100644 (file)
index 0000000..534c7c3
--- /dev/null
@@ -0,0 +1,2176 @@
+// Created on: 2013-08-27
+// Created by: Denis BOGOLEPOV
+// Copyright (c) 2013 OPEN CASCADE SAS
+//
+// The content of this file is subject to the Open CASCADE Technology Public
+// License Version 6.5 (the "License"). You may not use the content of this file
+// except in compliance with the License. Please obtain a copy of the License
+// at http://www.opencascade.org and read it completely before using this file.
+//
+// The Initial Developer of the Original Code is Open CASCADE S.A.S., having its
+// main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France.
+//
+// The Original Code and all software distributed under the License is
+// distributed on an "AS IS" basis, without warranty of any kind, and the
+// Initial Developer hereby disclaims all such warranties, including without
+// limitation, any warranties of merchantability, fitness for a particular
+// purpose or non-infringement. Please see the License for the specific terms
+// and conditions governing the rights and limitations under the License.
+
+#ifdef HAVE_CONFIG_H
+  #include <config.h>
+#endif
+
+#ifdef HAVE_OPENCL
+
+#if defined(_WIN32)
+
+  #include <windows.h>
+  #include <wingdi.h>
+
+  #pragma comment (lib, "DelayImp.lib")
+  #pragma comment (lib, "OpenCL.lib")
+
+#elif defined(__APPLE__) && !defined(MACOSX_USE_GLX)
+  #include <OpenGL/CGLCurrent.h>
+#else
+  #include <GL/glx.h>
+#endif
+
+#include <OpenGl_Context.hxx>
+#include <OpenGl_Texture.hxx>
+#include <OpenGl_View.hxx>
+#include <OpenGl_Workspace.hxx>
+
+using namespace OpenGl_Raytrace;
+
+//! Use this macro to output ray-tracing debug info
+// #define RAY_TRACE_PRINT_INFO
+
+#ifdef RAY_TRACE_PRINT_INFO
+  #include <OSD_Timer.hxx>
+#endif
+
+//! OpenCL source of ray-tracing kernels.
+extern const char THE_RAY_TRACE_OPENCL_SOURCE[];
+
+// =======================================================================
+// function : MatVecMult
+// purpose  : Multiples 4x4 matrix by 4D vector
+// =======================================================================
+template< typename T >
+OpenGl_RTVec4f MatVecMult (const T m[16], const OpenGl_RTVec4f& v)
+{
+  return OpenGl_RTVec4f (
+    static_cast<float> (m[ 0] * v.x() + m[ 4] * v.y() +
+                        m[ 8] * v.z() + m[12] * v.w()),
+    static_cast<float> (m[ 1] * v.x() + m[ 5] * v.y() +
+                        m[ 9] * v.z() + m[13] * v.w()),
+    static_cast<float> (m[ 2] * v.x() + m[ 6] * v.y() +
+                        m[10] * v.z() + m[14] * v.w()),
+    static_cast<float> (m[ 3] * v.x() + m[ 7] * v.y() +
+                        m[11] * v.z() + m[15] * v.w()));
+}
+
+// =======================================================================
+// function : UpdateRaytraceEnvironmentMap
+// purpose  : Updates environment map for ray-tracing
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::UpdateRaytraceEnvironmentMap()
+{
+  if (myView.IsNull())
+    return Standard_False;
+
+  if (myViewModificationStatus == myView->ModificationState())
+    return Standard_True;
+
+  cl_int anError = CL_SUCCESS;
+
+  if (myRaytraceEnvironment != NULL)
+    clReleaseMemObject (myRaytraceEnvironment);
+
+  int aSizeX = 1;
+  int aSizeY = 1;
+
+  if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE)
+  {
+    aSizeX = (myView->TextureEnv()->SizeX() <= 0) ? 1 : myView->TextureEnv()->SizeX();
+    aSizeY = (myView->TextureEnv()->SizeY() <= 0) ? 1 : myView->TextureEnv()->SizeY();
+  }
+
+  cl_image_format aImageFormat;
+
+  aImageFormat.image_channel_order = CL_RGBA;
+  aImageFormat.image_channel_data_type = CL_FLOAT;
+
+  myRaytraceEnvironment = clCreateImage2D (myComputeContext, CL_MEM_READ_ONLY,
+                                           &aImageFormat, aSizeX, aSizeY, 0,
+                                           NULL, &anError);
+
+  cl_float* aPixelData = new cl_float[aSizeX * aSizeY * 4];
+
+  // Note: texture format is not compatible with OpenCL image
+  // (it's not possible to create image directly from texture)
+
+  if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE)
+  {
+    myView->TextureEnv()->Bind (GetGlContext());
+
+    glGetTexImage (GL_TEXTURE_2D,
+                   0,
+                   GL_RGBA,
+                   GL_FLOAT,
+                   aPixelData);
+
+    myView->TextureEnv()->Unbind (GetGlContext());
+  }
+  else
+  {
+    for (int aPixel = 0; aPixel < aSizeX * aSizeY * 4; ++aPixel)
+      aPixelData[aPixel] = 0.f;
+  }
+
+  size_t anImageOffset[] = { 0,
+                             0,
+                             0 };
+
+  size_t anImageRegion[] = { aSizeX,
+                             aSizeY,
+                             1 };
+
+  anError |= clEnqueueWriteImage (myRaytraceQueue, myRaytraceEnvironment, CL_TRUE,
+                                  anImageOffset, anImageRegion, 0, 0, aPixelData,
+                                  0, NULL, NULL);
+#ifdef RAY_TRACE_PRINT_INFO
+  if (anError != CL_SUCCESS)
+    std::cout << "Error! Failed to write environment map image!" << std::endl;
+#endif
+
+  delete[] aPixelData;
+
+  myViewModificationStatus = myView->ModificationState();
+
+  return (anError == CL_SUCCESS);
+}
+
+// =======================================================================
+// function : UpdateRaytraceGeometry
+// purpose  : Updates 3D scene geometry for ray tracing
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theCheck)
+{
+  if (myView.IsNull())
+    return Standard_False;
+
+  // Note: In 'check' mode the scene geometry is analyzed for modifications
+  // This is light-weight procedure performed for each frame
+
+  if (!theCheck)
+  {
+    myRaytraceSceneData.Clear();
+
+    myIsRaytraceDataValid = Standard_False;
+  }
+  else
+  {
+    if (myLayersModificationStatus != myView->LayerList().ModificationState())
+    {
+      return UpdateRaytraceGeometry (Standard_False);
+    }
+  }
+
+  float* aTransform (NULL);
+
+  // The set of processed structures (reflected to ray-tracing)
+  // This set is used to remove out-of-date records from the
+  // hash map of structures
+  std::set<const OpenGl_Structure*> anElements;
+
+  const OpenGl_LayerList& aList = myView->LayerList();
+
+  for (OpenGl_SequenceOfLayers::Iterator anLayerIt (aList.Layers()); anLayerIt.More(); anLayerIt.Next())
+  {
+    const OpenGl_PriorityList& aPriorityList = anLayerIt.Value();
+
+    if (aPriorityList.NbStructures() == 0)
+      continue;
+
+    const OpenGl_ArrayOfStructure& aStructArray = aPriorityList.ArrayOfStructures();
+
+    for (int anIndex = 0; anIndex < aStructArray.Length(); ++anIndex)
+    {
+      OpenGl_SequenceOfStructure::Iterator aStructIt;
+
+      for (aStructIt.Init (aStructArray (anIndex)); aStructIt.More(); aStructIt.Next())
+      {
+        const OpenGl_Structure* aStructure = aStructIt.Value();
+
+        if (theCheck)
+        {
+          if (CheckRaytraceStructure (aStructure))
+          {
+            return UpdateRaytraceGeometry (Standard_False);
+          }
+        }
+        else
+        {
+          if (!aStructure->IsRaytracable())
+            continue;
+
+          if (aStructure->Transformation()->mat != NULL)
+          {
+            if (aTransform == NULL)
+              aTransform = new float[16];
+
+            for (int i = 0; i < 4; ++i)
+              for (int j = 0; j < 4; ++j)
+              {
+                aTransform[j * 4 + i] = aStructure->Transformation()->mat[i][j];
+              }
+          }
+
+          AddRaytraceStructure (aStructure, aTransform, anElements);
+        }
+      }
+    }
+  }
+
+  if (!theCheck)
+  {
+    // Actualize the hash map of structures -- remove out-of-date records
+    std::map<const OpenGl_Structure*, Standard_Size>::iterator anIter = myStructureStates.begin();
+
+    while (anIter != myStructureStates.end())
+    {
+      if (anElements.find (anIter->first) == anElements.end())
+      {
+        myStructureStates.erase (anIter++);
+      }
+      else
+      {
+        ++anIter;
+      }
+    }
+
+    // Actualize OpenGL layer list state
+    myLayersModificationStatus = myView->LayerList().ModificationState();
+
+
+#ifdef RAY_TRACE_PRINT_INFO
+    OSD_Timer aTimer;
+    aTimer.Start();
+#endif
+
+    myBVHBuilder.Build (myRaytraceSceneData);
+
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << " Build time: " << aTimer.ElapsedTime() << " for "
+                        << myRaytraceSceneData.Triangles.size() / 1000 << "K triangles" << std::endl;
+#endif
+
+    const float aScaleFactor = 1.5f;
+
+    myRaytraceSceneRadius = aScaleFactor *
+      Max ( Max (fabsf (myRaytraceSceneData.AABB.CornerMin().x()),
+            Max (fabsf (myRaytraceSceneData.AABB.CornerMin().y()),
+                 fabsf (myRaytraceSceneData.AABB.CornerMin().z()))),
+            Max (fabsf (myRaytraceSceneData.AABB.CornerMax().x()),
+            Max (fabsf (myRaytraceSceneData.AABB.CornerMax().y()),
+                 fabsf (myRaytraceSceneData.AABB.CornerMax().z()))) );
+
+    myRaytraceSceneEpsilon = Max (1e-4f, myRaytraceSceneRadius * 1e-4f);
+
+    return WriteRaytraceSceneToDevice();
+  }
+
+  delete [] aTransform;
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : CheckRaytraceStructure
+// purpose  : Adds OpenGL structure to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::CheckRaytraceStructure (const OpenGl_Structure* theStructure)
+{
+  if (!theStructure->IsRaytracable())
+  {
+    // Checks to see if all ray-tracable elements were
+    // removed from the structure
+    if (theStructure->ModificationState() > 0)
+    {
+      theStructure->ResetModificationState();
+      return Standard_True;
+    }
+
+    return Standard_False;
+  }
+
+  std::map<const OpenGl_Structure*, Standard_Size>::iterator aStructState = myStructureStates.find (theStructure);
+
+  if (aStructState != myStructureStates.end())
+    return aStructState->second != theStructure->ModificationState();
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : CreateMaterial
+// purpose  : Creates ray-tracing material properties
+// =======================================================================
+void CreateMaterial (const OPENGL_SURF_PROP& theProp, OpenGl_RaytraceMaterial& theMaterial)
+{
+  theMaterial.Ambient = OpenGl_RTVec4f (theProp.ambcol.rgb[0] * theProp.amb,
+                                        theProp.ambcol.rgb[1] * theProp.amb,
+                                        theProp.ambcol.rgb[2] * theProp.amb,
+                                        1.f);
+
+  theMaterial.Diffuse = OpenGl_RTVec4f (theProp.difcol.rgb[0] * theProp.diff,
+                                        theProp.difcol.rgb[1] * theProp.diff,
+                                        theProp.difcol.rgb[2] * theProp.diff,
+                                        1.f);
+
+  theMaterial.Specular = OpenGl_RTVec4f (theProp.speccol.rgb[0] * theProp.spec,
+                                         theProp.speccol.rgb[1] * theProp.spec,
+                                         theProp.speccol.rgb[2] * theProp.spec,
+                                         theProp.shine);
+
+  theMaterial.Emission = OpenGl_RTVec4f (theProp.emscol.rgb[0] * theProp.emsv,
+                                         theProp.emscol.rgb[1] * theProp.emsv,
+                                         theProp.emscol.rgb[2] * theProp.emsv,
+                                         1.f);
+
+  // Note: Here we use sub-linear transparency function
+  // to produce realistic-looking transparency effect
+  theMaterial.Transparency = OpenGl_RTVec4f (powf (theProp.trans, 0.75f),
+                                             1.f - theProp.trans,
+                                             1.f,
+                                             1.f);
+
+  const float aMaxRefl = Max (theMaterial.Diffuse.x() + theMaterial.Specular.x(),
+                         Max (theMaterial.Diffuse.y() + theMaterial.Specular.y(),
+                              theMaterial.Diffuse.z() + theMaterial.Specular.z()));
+
+  const float aReflectionScale = 0.75f / aMaxRefl;
+
+  theMaterial.Reflection = OpenGl_RTVec4f (theProp.speccol.rgb[0] * theProp.spec,
+                                           theProp.speccol.rgb[1] * theProp.spec,
+                                           theProp.speccol.rgb[2] * theProp.spec,
+                                           0.f) * aReflectionScale;
+}
+
+// =======================================================================
+// function : AddRaytraceStructure
+// purpose  : Adds OpenGL structure to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure*            theStructure,
+                                                         const float*                       theTransform,
+                                                         std::set<const OpenGl_Structure*>& theElements)
+{
+#ifdef RAY_TRACE_PRINT_INFO
+  std::cout << "Add Structure" << std::endl;
+#endif
+
+  theElements.insert (theStructure);
+
+  if (!theStructure->IsVisible())
+  {
+    myStructureStates[theStructure] = theStructure->ModificationState();
+    return Standard_True;
+  }
+
+  // Get structure material
+  int aStructMatID = -1;
+
+  if (theStructure->AspectFace() != NULL)
+  {
+    aStructMatID = static_cast<int> (myRaytraceSceneData.Materials.size());
+
+    OpenGl_RaytraceMaterial aStructMaterial;
+    CreateMaterial (theStructure->AspectFace()->IntFront(), aStructMaterial);
+
+    myRaytraceSceneData.Materials.push_back (aStructMaterial);
+  }
+
+  OpenGl_ListOfGroup::Iterator anItg (theStructure->Groups());
+
+  while (anItg.More())
+  {
+    // Get group material
+    int aGroupMatID = -1;
+
+    if (anItg.Value()->AspectFace() != NULL)
+    {
+      aGroupMatID = static_cast<int> (myRaytraceSceneData.Materials.size());
+
+      OpenGl_RaytraceMaterial aGroupMaterial;
+      CreateMaterial (anItg.Value()->AspectFace()->IntFront(), aGroupMaterial);
+
+      myRaytraceSceneData.Materials.push_back (aGroupMaterial);
+    }
+
+    int aMatID = aGroupMatID < 0 ? aStructMatID : aGroupMatID;
+
+    if (aStructMatID < 0 && aGroupMatID < 0)
+    {
+      aMatID = static_cast<int> (myRaytraceSceneData.Materials.size());
+
+      myRaytraceSceneData.Materials.push_back (OpenGl_RaytraceMaterial());
+    }
+
+    // Add OpenGL elements from group (only arrays of primitives)
+    for (const OpenGl_ElementNode* aNode = anItg.Value()->FirstNode(); aNode != NULL; aNode = aNode->next)
+    {
+      if (TelNil == aNode->type)
+      {
+        OpenGl_AspectFace* anAspect = dynamic_cast<OpenGl_AspectFace*> (aNode->elem);
+
+        if (anAspect != NULL)
+        {
+          aMatID = static_cast<int> (myRaytraceSceneData.Materials.size());
+
+          OpenGl_RaytraceMaterial aMaterial;
+          CreateMaterial (anAspect->IntFront(), aMaterial);
+
+          myRaytraceSceneData.Materials.push_back (aMaterial);
+        }
+      }
+      else if (TelParray == aNode->type)
+      {
+        OpenGl_PrimitiveArray* aPrimArray = dynamic_cast<OpenGl_PrimitiveArray*> (aNode->elem);
+
+        if (aPrimArray != NULL)
+        {
+          AddRaytracePrimitiveArray (aPrimArray->PArray(), aMatID, theTransform);
+        }
+      }
+    }
+
+    anItg.Next();
+  }
+
+  float* aTransform (NULL);
+
+  // Process all connected OpenGL structures
+  OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures());
+
+  while (anIts.More())
+  {
+    if (anIts.Value()->Transformation()->mat != NULL)
+    {
+      float* aTransform = new float[16];
+
+      for (int i = 0; i < 4; ++i)
+        for (int j = 0; j < 4; ++j)
+        {
+          aTransform[j * 4 + i] =
+            anIts.Value()->Transformation()->mat[i][j];
+        }
+    }
+
+    if (anIts.Value()->IsRaytracable())
+      AddRaytraceStructure (anIts.Value(), aTransform != NULL ? aTransform : theTransform, theElements);
+
+    anIts.Next();
+  }
+
+  delete[] aTransform;
+
+  myStructureStates[theStructure] = theStructure->ModificationState();
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : AddRaytracePrimitiveArray
+// purpose  : Adds OpenGL primitive array to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::AddRaytracePrimitiveArray (const CALL_DEF_PARRAY* theArray,
+                                                              int                    theMatID,
+                                                              const float*           theTransform)
+{
+  if (theArray->type != TelPolygonsArrayType &&
+      theArray->type != TelTrianglesArrayType &&
+      theArray->type != TelQuadranglesArrayType &&
+      theArray->type != TelTriangleFansArrayType &&
+      theArray->type != TelTriangleStripsArrayType &&
+      theArray->type != TelQuadrangleStripsArrayType)
+  {
+    return Standard_True;
+  }
+
+  if (theArray->vertices == NULL)
+    return Standard_False;
+
+#ifdef RAY_TRACE_PRINT_INFO
+  switch (theArray->type)
+  {
+    case TelPolygonsArrayType:
+      std::cout << "\tTelPolygonsArrayType" << std::endl; break;
+    case TelTrianglesArrayType:
+      std::cout << "\tTelTrianglesArrayType" << std::endl; break;
+    case TelQuadranglesArrayType:
+      std::cout << "\tTelQuadranglesArrayType" << std::endl; break;
+    case TelTriangleFansArrayType:
+      std::cout << "\tTelTriangleFansArrayType" << std::endl; break;
+    case TelTriangleStripsArrayType:
+      std::cout << "\tTelTriangleStripsArrayType" << std::endl; break;
+    case TelQuadrangleStripsArrayType:
+      std::cout << "\tTelQuadrangleStripsArrayType" << std::endl; break;
+  }
+#endif
+
+  // Simple optimization to eliminate possible memory allocations
+  // during processing of the primitive array vertices
+  myRaytraceSceneData.Vertices.reserve (
+    myRaytraceSceneData.Vertices.size() + theArray->num_vertexs);
+
+  const int aFirstVert = static_cast<int> (myRaytraceSceneData.Vertices.size());
+
+  for (int aVert = 0; aVert < theArray->num_vertexs; ++aVert)
+  {
+    OpenGl_RTVec4f aVertex (theArray->vertices[aVert].xyz[0],
+                            theArray->vertices[aVert].xyz[1],
+                            theArray->vertices[aVert].xyz[2],
+                            1.f);
+
+    if (theTransform)
+      aVertex = MatVecMult (theTransform, aVertex);
+
+    myRaytraceSceneData.Vertices.push_back (aVertex);
+
+    myRaytraceSceneData.AABB.Add (aVertex);
+  }
+
+  myRaytraceSceneData.Normals.reserve (
+    myRaytraceSceneData.Normals.size() + theArray->num_vertexs);
+
+  for (int aNorm = 0; aNorm < theArray->num_vertexs; ++aNorm)
+  {
+    OpenGl_RTVec4f aNormal;
+
+    // Note: In case of absence of normals, the visualizer
+    // will use generated geometric normals
+
+    if (theArray->vnormals != NULL)
+    {
+      aNormal = OpenGl_RTVec4f (theArray->vnormals[aNorm].xyz[0],
+                                theArray->vnormals[aNorm].xyz[1],
+                                theArray->vnormals[aNorm].xyz[2],
+                                0.f);
+
+      if (theTransform)
+        aNormal = MatVecMult (theTransform, aNormal);
+    }
+
+    myRaytraceSceneData.Normals.push_back (aNormal);
+  }
+
+  if (theArray->num_bounds > 0)
+  {
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "\tNumber of bounds = " << theArray->num_bounds << std::endl;
+#endif
+
+    int aVertOffset = 0;
+
+    for (int aBound = 0; aBound < theArray->num_bounds; ++aBound)
+    {
+      const int aVertNum = theArray->bounds[aBound];
+
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "\tAdd indices from bound " << aBound << ": " <<
+                                    aVertOffset << ", " << aVertNum << std::endl;
+#endif
+
+      if (!AddRaytraceVertexIndices (theArray, aFirstVert, aVertOffset, aVertNum, theMatID))
+      {
+        return Standard_False;
+      }
+
+      aVertOffset += aVertNum;
+    }
+  }
+  else
+  {
+    const int aVertNum = theArray->num_edges > 0 ? theArray->num_edges : theArray->num_vertexs;
+
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "\tAdd indices: " << aVertNum << std::endl;
+#endif
+
+    return AddRaytraceVertexIndices (theArray, aFirstVert, 0, aVertNum, theMatID);
+  }
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : AddRaytraceVertexIndices
+// purpose  : Adds vertex indices to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::AddRaytraceVertexIndices (const CALL_DEF_PARRAY* theArray,
+                                                             int                    theFirstVert,
+                                                             int                    theVertOffset,
+                                                             int                    theVertNum,
+                                                             int                    theMatID)
+{
+  myRaytraceSceneData.Triangles.reserve (myRaytraceSceneData.Triangles.size() + theVertNum);
+
+  switch (theArray->type)
+  {
+    case TelTrianglesArrayType:
+      return AddRaytraceTriangleArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
+
+    case TelQuadranglesArrayType:
+      return AddRaytraceQuadrangleArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
+
+    case TelTriangleFansArrayType:
+      return AddRaytraceTriangleFanArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
+
+    case TelTriangleStripsArrayType:
+      return AddRaytraceTriangleStripArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
+
+    case TelQuadrangleStripsArrayType:
+      return AddRaytraceQuadrangleStripArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
+
+    case TelPolygonsArrayType:
+      return AddRaytracePolygonArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
+  }
+
+  return Standard_False;
+}
+
+// =======================================================================
+// function : AddRaytraceTriangleArray
+// purpose  : Adds OpenGL triangle array to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleArray (const CALL_DEF_PARRAY* theArray,
+                                                             int                    theFirstVert,
+                                                             int                    theVertOffset,
+                                                             int                    theVertNum,
+                                                             int                    theMatID)
+{
+  if (theVertNum < 3)
+    return Standard_True;
+
+  if (theArray->num_edges > 0)
+  {
+    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; aVert += 3)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[aVert + 0],
+                                                               theFirstVert + theArray->edges[aVert + 1],
+                                                               theFirstVert + theArray->edges[aVert + 2],
+                                                               theMatID));
+    }
+  }
+  else
+  {
+    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; aVert += 3)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0,
+                                                               theFirstVert + aVert + 1,
+                                                               theFirstVert + aVert + 2,
+                                                               theMatID));
+    }
+  }
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : AddRaytraceTriangleFanArray
+// purpose  : Adds OpenGL triangle fan array to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleFanArray (const CALL_DEF_PARRAY* theArray,
+                                                                int                     theFirstVert,
+                                                                int                     theVertOffset,
+                                                                int                     theVertNum,
+                                                                int                     theMatID)
+{
+  if (theVertNum < 3)
+    return Standard_True;
+
+  if (theArray->num_edges > 0)
+  {
+    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[theVertOffset],
+                                                               theFirstVert + theArray->edges[aVert + 1],
+                                                               theFirstVert + theArray->edges[aVert + 2],
+                                                               theMatID));
+    }
+  }
+  else
+  {
+    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset,
+                                                               theFirstVert + aVert + 1,
+                                                               theFirstVert + aVert + 2,
+                                                               theMatID));
+    }
+  }
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : AddRaytraceTriangleStripArray
+// purpose  : Adds OpenGL triangle strip array to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleStripArray (const CALL_DEF_PARRAY* theArray,
+                                                                  int                    theFirstVert,
+                                                                  int                    theVertOffset,
+                                                                  int                    theVertNum,
+                                                                  int                    theMatID)
+{
+  if (theVertNum < 3)
+    return Standard_True;
+
+  if (theArray->num_edges > 0)
+  {
+    myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
+                                      theFirstVert + theArray->edges[theVertOffset + 0],
+                                      theFirstVert + theArray->edges[theVertOffset + 1],
+                                      theFirstVert + theArray->edges[theVertOffset + 2],
+                                      theMatID));
+
+    for (int aVert = theVertOffset + 1, aTriNum = 1; aVert < theVertOffset + theVertNum - 2; ++aVert, ++aTriNum)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
+                                      theFirstVert + theArray->edges[aVert + (aTriNum % 2) ? 1 : 0],
+                                      theFirstVert + theArray->edges[aVert + (aTriNum % 2) ? 0 : 1],
+                                      theFirstVert + theArray->edges[aVert + 2],
+                                      theMatID));
+    }
+  }
+  else
+  {
+    myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset + 0,
+                                                             theFirstVert + theVertOffset + 1,
+                                                             theFirstVert + theVertOffset + 2,
+                                                             theMatID));
+
+    for (int aVert = theVertOffset + 1, aTriNum = 1; aVert < theVertOffset + theVertNum - 2; ++aVert, ++aTriNum)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + ( aTriNum % 2 ) ? 1 : 0,
+                                                               theFirstVert + aVert + ( aTriNum % 2 ) ? 0 : 1,
+                                                               theFirstVert + aVert + 2,
+                                                               theMatID));
+    }
+  }
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : AddRaytraceQuadrangleArray
+// purpose  : Adds OpenGL quad array to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleArray (const CALL_DEF_PARRAY* theArray,
+                                                               int                    theFirstVert,
+                                                               int                    theVertOffset,
+                                                               int                    theVertNum,
+                                                               int                    theMatID)
+{
+  if (theVertNum < 4)
+    return Standard_True;
+
+  if (theArray->num_edges > 0)
+  {
+    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 3; aVert += 4)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[aVert + 0],
+                                                               theFirstVert + theArray->edges[aVert + 1],
+                                                               theFirstVert + theArray->edges[aVert + 2],
+                                                               theMatID));
+
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[aVert + 0],
+                                                               theFirstVert + theArray->edges[aVert + 2],
+                                                               theFirstVert + theArray->edges[aVert + 3],
+                                                               theMatID));
+    }
+  }
+  else
+  {
+    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 3; aVert += 4)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0,
+                                                               theFirstVert + aVert + 1,
+                                                               theFirstVert + aVert + 2,
+                                                               theMatID));
+
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0,
+                                                               theFirstVert + aVert + 2,
+                                                               theFirstVert + aVert + 3,
+                                                               theMatID));
+    }
+  }
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : AddRaytraceQuadrangleStripArray
+// purpose  : Adds OpenGL quad strip array to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleStripArray (const CALL_DEF_PARRAY* theArray,
+                                                                    int                    theFirstVert,
+                                                                    int                    theVertOffset,
+                                                                    int                    theVertNum,
+                                                                    int                    theMatID)
+{
+  if (theVertNum < 4)
+    return Standard_True;
+
+  if (theArray->num_edges > 0)
+  {
+    myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
+                                theFirstVert + theArray->edges[theVertOffset + 0],
+                                theFirstVert + theArray->edges[theVertOffset + 1],
+                                theFirstVert + theArray->edges[theVertOffset + 2],
+                                theMatID));
+
+    myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
+                                theFirstVert + theArray->edges[theVertOffset + 1],
+                                theFirstVert + theArray->edges[theVertOffset + 3],
+                                theFirstVert + theArray->edges[theVertOffset + 2],
+                                theMatID));
+
+    for (int aVert = theVertOffset + 2; aVert < theVertOffset + theVertNum - 3; aVert += 2)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
+                                  theFirstVert + theArray->edges[aVert + 0],
+                                  theFirstVert + theArray->edges[aVert + 1],
+                                  theFirstVert + theArray->edges[aVert + 2],
+                                  theMatID));
+
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
+                                  theFirstVert + theArray->edges[aVert + 1],
+                                  theFirstVert + theArray->edges[aVert + 3],
+                                  theFirstVert + theArray->edges[aVert + 2],
+                                  theMatID));
+    }
+  }
+  else
+  {
+    myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + 0,
+                                                             theFirstVert + 1,
+                                                             theFirstVert + 2,
+                                                             theMatID));
+
+    myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + 1,
+                                                             theFirstVert + 3,
+                                                             theFirstVert + 2,
+                                                             theMatID));
+
+    for (int aVert = theVertOffset + 2; aVert < theVertOffset + theVertNum - 3; aVert += 2)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0,
+                                                               theFirstVert + aVert + 1,
+                                                               theFirstVert + aVert + 2,
+                                                               theMatID));
+
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 1,
+                                                               theFirstVert + aVert + 3,
+                                                               theFirstVert + aVert + 2,
+                                                               theMatID));
+    }
+  }
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : AddRaytracePolygonArray
+// purpose  : Adds OpenGL polygon array to ray-traced scene geometry
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::AddRaytracePolygonArray (const CALL_DEF_PARRAY* theArray,
+                                                            int                    theFirstVert,
+                                                            int                    theVertOffset,
+                                                            int                    theVertNum,
+                                                            int                    theMatID)
+{
+  if (theArray->num_vertexs < 3)
+    return Standard_True;
+
+  if (theArray->edges != NULL)
+  {
+    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[theVertOffset],
+                                                               theFirstVert + theArray->edges[aVert + 1],
+                                                               theFirstVert + theArray->edges[aVert + 2],
+                                                               theMatID));
+    }
+  }
+  else
+  {
+    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+    {
+      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset,
+                                                               theFirstVert + aVert + 1,
+                                                               theFirstVert + aVert + 2,
+                                                               theMatID));
+    }
+  }
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : UpdateRaytraceLightSources
+// purpose  : Updates 3D scene light sources for ray-tracing
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble theInvModelView[16])
+{
+  myRaytraceSceneData.LightSources.clear();
+
+  OpenGl_ListOfLight::Iterator anItl (myView->LightList());
+
+  OpenGl_RTVec4f aAmbient (0.f, 0.f, 0.f, 0.f);
+
+  for (; anItl.More(); anItl.Next())
+  {
+    const OpenGl_Light &aLight = anItl.Value();
+
+    if (aLight.type == TLightAmbient)
+    {
+      aAmbient += OpenGl_RTVec4f (aLight.col.rgb[0],
+                                  aLight.col.rgb[1],
+                                  aLight.col.rgb[2],
+                                  0.f);
+      continue;
+    }
+
+    OpenGl_RTVec4f aDiffuse (aLight.col.rgb[0],
+                             aLight.col.rgb[1],
+                             aLight.col.rgb[2],
+                             1.f);
+
+    OpenGl_RTVec4f aPosition (-aLight.dir[0],
+                              -aLight.dir[1],
+                              -aLight.dir[2],
+                              0.f);
+
+    if (aLight.type != TLightDirectional)
+    {
+      aPosition = OpenGl_RTVec4f (aLight.pos[0],
+                                  aLight.pos[1],
+                                  aLight.pos[2],
+                                  1.f);
+    }
+
+    if (aLight.HeadLight)
+    {
+      aPosition = MatVecMult (theInvModelView, aPosition);
+    }
+
+    myRaytraceSceneData.LightSources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition));
+  }
+
+  if (myRaytraceSceneData.LightSources.size() > 0)
+  {
+    myRaytraceSceneData.LightSources.front().Ambient += aAmbient;
+  }
+  else
+  {
+    myRaytraceSceneData.LightSources.push_back (OpenGl_RaytraceLight (OpenGl_RTVec4f (aAmbient.xyz(), -1.0f)));
+  }
+
+  cl_int anError = CL_SUCCESS;
+
+  if (myRaytraceLightSourceBuffer != NULL)
+    clReleaseMemObject (myRaytraceLightSourceBuffer);
+
+  const size_t myLightBufferSize = myRaytraceSceneData.LightSources.size() > 0
+                                 ? myRaytraceSceneData.LightSources.size()
+                                 : 1;
+
+  myRaytraceLightSourceBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
+                                                myLightBufferSize * sizeof(OpenGl_RaytraceLight),
+                                                NULL, &anError);
+
+  if (myRaytraceSceneData.LightSources.size() > 0)
+  {
+    const void* aDataPtr = myRaytraceSceneData.LightSources.front().Packed();
+    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0,
+                                     myLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr,
+                                     0, NULL, NULL);
+  }
+
+#ifdef RAY_TRACE_PRINT_INFO
+  if (anError != CL_SUCCESS)
+  {
+    std::cout << "Error! Failed to set light sources!";
+
+    return Standard_False;
+  }
+#endif
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : CheckOpenCL
+// purpose  : Checks OpenCL dynamic library availability
+// =======================================================================
+Standard_Boolean CheckOpenCL()
+{
+#if defined ( _WIN32 )
+
+  __try
+  {
+    cl_uint aNbPlatforms;
+    clGetPlatformIDs (0, NULL, &aNbPlatforms);
+  }
+  __except (EXCEPTION_EXECUTE_HANDLER)
+  {
+    return Standard_False;
+  }
+
+#endif
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : InitOpenCL
+// purpose  : Initializes OpenCL objects
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::InitOpenCL()
+{
+  if (myComputeInitStatus != OpenGl_CLIS_NONE)
+  {
+    return myComputeInitStatus == OpenGl_CLIS_INIT;
+  }
+
+  if (!CheckOpenCL())
+  {
+    myComputeInitStatus = OpenGl_CLIS_FAIL; // fail to load OpenCL library
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                              GL_DEBUG_TYPE_ERROR_ARB,
+                              0,
+                              GL_DEBUG_SEVERITY_HIGH_ARB,
+                              "Failed to load OpenCL dynamic library!");
+    return Standard_False;
+  }
+
+  // Obtain the list of platforms available
+  cl_uint aNbPlatforms = 0;
+  cl_int  anError = clGetPlatformIDs (0, NULL, &aNbPlatforms);
+  cl_platform_id* aPlatforms = (cl_platform_id* )alloca (aNbPlatforms * sizeof(cl_platform_id));
+  anError |= clGetPlatformIDs (aNbPlatforms, aPlatforms, NULL);
+  if (anError != CL_SUCCESS
+   || aNbPlatforms == 0)
+  {
+    myComputeInitStatus = OpenGl_CLIS_FAIL;
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                              GL_DEBUG_TYPE_ERROR_ARB,
+                              0,
+                              GL_DEBUG_SEVERITY_HIGH_ARB,
+                              "No any OpenCL platform installed!");
+    return Standard_False;
+  }
+
+  // Note: We try to find NVIDIA or AMD platforms with GPU devices!
+  cl_platform_id aPrefPlatform = NULL;
+  for (cl_uint aPlatIter = 0; aPlatIter < aNbPlatforms; ++aPlatIter)
+  {
+    char aName[256];
+    anError = clGetPlatformInfo (aPlatforms[aPlatIter], CL_PLATFORM_NAME,
+                                 sizeof(aName), aName, NULL);
+    if (anError != CL_SUCCESS)
+    {
+      continue;
+    }
+
+    if (strncmp (aName, "NVIDIA", strlen ("NVIDIA")) == 0)
+    {
+      aPrefPlatform = aPlatforms[aPlatIter];
+
+      // Use optimizations for NVIDIA GPUs
+      myIsAmdComputePlatform = Standard_False;
+    }
+    else if (strncmp (aName, "AMD", strlen ("AMD")) == 0)
+    {
+      aPrefPlatform = (aPrefPlatform == NULL)
+                    ? aPlatforms[aPlatIter]
+                    : aPrefPlatform;
+
+      // Use optimizations for ATI/AMD platform
+      myIsAmdComputePlatform = Standard_True;
+    }
+  }
+
+  if (aPrefPlatform == NULL)
+  {
+    aPrefPlatform = aPlatforms[0];
+  }
+
+  // Obtain the list of devices available in the selected platform
+  cl_uint aNbDevices = 0;
+  anError = clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU,
+                            0, NULL, &aNbDevices);
+
+  cl_device_id* aDevices = (cl_device_id* )alloca (aNbDevices * sizeof(cl_device_id));
+  anError |= clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU,
+                             aNbDevices, aDevices, NULL);
+  if (anError != CL_SUCCESS)
+  {
+    myComputeInitStatus = OpenGl_CLIS_FAIL;
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                              GL_DEBUG_TYPE_ERROR_ARB,
+                              0,
+                              GL_DEBUG_SEVERITY_HIGH_ARB,
+                              "Failed to get OpenCL GPU device!");
+    return Standard_False;
+  }
+
+  // Note: Simply get first available GPU
+  cl_device_id aDevice = aDevices[0];
+
+  // detect old contexts
+  char aVerClStr[256];
+  clGetDeviceInfo (aDevice, CL_DEVICE_VERSION,
+                   sizeof(aVerClStr), aVerClStr, NULL);
+  aVerClStr[strlen ("OpenCL 1.0")] = '\0';
+  const bool isVer10 = strncmp (aVerClStr, "OpenCL 1.0", strlen ("OpenCL 1.0")) == 0;
+
+  // Create OpenCL context
+  cl_context_properties aCtxProp[] =
+  {
+  #if defined(__APPLE__) && !defined(MACOSX_USE_GLX)
+    CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE,
+    (cl_context_properties )CGLGetShareGroup (CGLGetCurrentContext()),
+  #elif defined(_WIN32)
+    CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform,
+    CL_GL_CONTEXT_KHR,   (cl_context_properties )wglGetCurrentContext(),
+    CL_WGL_HDC_KHR,      (cl_context_properties )wglGetCurrentDC(),
+  #else
+    CL_GL_CONTEXT_KHR,   (cl_context_properties )glXGetCurrentContext(),
+    CL_GLX_DISPLAY_KHR,  (cl_context_properties )glXGetCurrentDisplay(),
+    CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform,
+  #endif
+    0
+  };
+
+  myComputeContext = clCreateContext (aCtxProp,
+                                    #if defined(__APPLE__) && !defined(MACOSX_USE_GLX)
+                                      0, NULL, // device will be taken from GL context
+                                    #else
+                                      1, &aDevice,
+                                    #endif
+                                      NULL, NULL, &anError);
+  if (anError != CL_SUCCESS)
+  {
+    myComputeInitStatus = OpenGl_CLIS_FAIL;
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                              GL_DEBUG_TYPE_ERROR_ARB,
+                              0,
+                              GL_DEBUG_SEVERITY_HIGH_ARB,
+                              "Failed to initialize OpenCL context!");
+    return Standard_False;
+  }
+
+  // Create OpenCL program
+  const char* aSources[] =
+  {
+    isVer10 ? "#define M_PI_F ( float )( 3.14159265359f )\n" : "",
+    THE_RAY_TRACE_OPENCL_SOURCE
+  };
+  myRaytraceProgram = clCreateProgramWithSource (myComputeContext, 2,
+                                                 aSources, NULL, &anError);
+  if (anError != CL_SUCCESS)
+  {
+    myComputeInitStatus = OpenGl_CLIS_FAIL;
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                              GL_DEBUG_TYPE_ERROR_ARB,
+                              0,
+                              GL_DEBUG_SEVERITY_HIGH_ARB,
+                              "Failed to create OpenCL ray-tracing program!");
+    return Standard_False;
+  }
+
+  anError = clBuildProgram (myRaytraceProgram, 0,
+                            NULL, NULL, NULL, NULL);
+  {
+    // Fetch build log
+    size_t aLogLen = 0;
+    cl_int aResult = clGetProgramBuildInfo (myRaytraceProgram, aDevice,
+                                            CL_PROGRAM_BUILD_LOG, 0, NULL, &aLogLen);
+
+    char* aBuildLog = (char* )alloca (aLogLen);
+    aResult |= clGetProgramBuildInfo (myRaytraceProgram, aDevice,
+                                      CL_PROGRAM_BUILD_LOG, aLogLen, aBuildLog, NULL);
+    if (aResult == CL_SUCCESS)
+    {
+      if (anError != CL_SUCCESS)
+      {
+        myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                                  GL_DEBUG_TYPE_ERROR_ARB,
+                                  0,
+                                  GL_DEBUG_SEVERITY_HIGH_ARB,
+                                  aBuildLog);
+      }
+      else
+      {
+      #ifdef RAY_TRACE_PRINT_INFO
+        std::cout << aBuildLog << std::endl;
+      #endif
+      }
+    }
+  }
+
+  if (anError != CL_SUCCESS)
+  {
+    return Standard_False;
+  }
+
+  // Create OpenCL ray tracing kernels
+  myRaytraceRenderKernel = clCreateKernel (myRaytraceProgram, "Main",            &anError);
+  if (anError != CL_SUCCESS)
+  {
+    myComputeInitStatus = OpenGl_CLIS_FAIL;
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                              GL_DEBUG_TYPE_ERROR_ARB,
+                              0,
+                              GL_DEBUG_SEVERITY_HIGH_ARB,
+                              "Failed to create OpenCL ray-tracing kernel!");
+    return Standard_False;
+  }
+
+  myRaytraceSmoothKernel = clCreateKernel (myRaytraceProgram, "MainAntialiased", &anError);
+  if (anError != CL_SUCCESS)
+  {
+    myComputeInitStatus = OpenGl_CLIS_FAIL;
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                              GL_DEBUG_TYPE_ERROR_ARB,
+                              0,
+                              GL_DEBUG_SEVERITY_HIGH_ARB,
+                              "Failed to create OpenCL ray-tracing kernel!");
+    return Standard_False;
+  }
+
+  // Create OpenCL command queue
+  // Note: For profiling set CL_QUEUE_PROFILING_ENABLE
+  cl_command_queue_properties aProps = CL_QUEUE_PROFILING_ENABLE;
+
+  myRaytraceQueue = clCreateCommandQueue (myComputeContext, aDevice, aProps, &anError);
+  if (anError != CL_SUCCESS)
+  {
+    myComputeInitStatus = OpenGl_CLIS_FAIL;
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                              GL_DEBUG_TYPE_ERROR_ARB,
+                              0,
+                              GL_DEBUG_SEVERITY_HIGH_ARB,
+                              "Failed to create OpenCL command queue!");
+
+    return Standard_False;
+  }
+
+  myComputeInitStatus = OpenGl_CLIS_INIT; // initialized in normal way
+  return Standard_True;
+}
+
+// =======================================================================
+// function : GetOpenClDeviceInfo
+// purpose  : Returns information about device used for computations
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::GetOpenClDeviceInfo (NCollection_DataMap<TCollection_AsciiString,
+                                                                            TCollection_AsciiString>& theInfo) const
+{
+  theInfo.Clear();
+  if (myComputeContext == NULL)
+  {
+    return Standard_False;
+  }
+
+  size_t aDevicesSize = 0;
+  cl_int anError = clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, 0, NULL, &aDevicesSize);
+  cl_device_id* aDevices = (cl_device_id* )alloca (aDevicesSize);
+  anError |= clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, aDevicesSize, aDevices, NULL);
+  if (anError != CL_SUCCESS)
+  {
+    return Standard_False;
+  }
+
+  char aDeviceName[256];
+  anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_NAME, sizeof(aDeviceName), aDeviceName, NULL);
+  theInfo.Bind ("Name", aDeviceName);
+
+  char aDeviceVendor[256];
+  anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_VENDOR, sizeof(aDeviceVendor), aDeviceVendor, NULL);
+  theInfo.Bind ("Vendor", aDeviceVendor);
+
+  cl_device_type aDeviceType;
+  anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_TYPE, sizeof(aDeviceType), &aDeviceType, NULL);
+  theInfo.Bind ("Type", aDeviceType == CL_DEVICE_TYPE_GPU ? "GPU" : "CPU");
+  return Standard_True;
+}
+
+// =======================================================================
+// function : ReleaseOpenCL
+// purpose  : Releases resources of OpenCL objects
+// =======================================================================
+void OpenGl_Workspace::ReleaseOpenCL()
+{
+  clReleaseKernel (myRaytraceRenderKernel);
+  clReleaseKernel (myRaytraceSmoothKernel);
+
+  clReleaseProgram (myRaytraceProgram);
+  clReleaseCommandQueue (myRaytraceQueue);
+
+  clReleaseMemObject (myRaytraceOutputImage);
+  clReleaseMemObject (myRaytraceEnvironment);
+  clReleaseMemObject (myRaytraceOutputImageSmooth);
+
+  clReleaseMemObject (myRaytraceVertexBuffer);
+  clReleaseMemObject (myRaytraceNormalBuffer);
+  clReleaseMemObject (myRaytraceTriangleBuffer);
+
+  clReleaseMemObject (myRaytraceMaterialBuffer);
+  clReleaseMemObject (myRaytraceLightSourceBuffer);
+
+  clReleaseMemObject (myRaytraceNodeMinPointBuffer);
+  clReleaseMemObject (myRaytraceNodeMaxPointBuffer);
+  clReleaseMemObject (myRaytraceNodeDataRcrdBuffer);
+
+  clReleaseContext (myComputeContext);
+
+  if (glIsTexture (*myRaytraceOutputTexture))
+    glDeleteTextures (2, myRaytraceOutputTexture);
+}
+
+// =======================================================================
+// function : ResizeRaytraceOutputBuffer
+// purpose  : Resizes OpenCL output image
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::ResizeRaytraceOutputBuffer (const cl_int theSizeX,
+                                                               const cl_int theSizeY)
+{
+  if (myComputeContext == NULL)
+  {
+    return Standard_False;
+  }
+
+  bool toResize = true;
+  GLint aSizeX = -1;
+  GLint aSizeY = -1;
+  if (*myRaytraceOutputTexture != 0)
+  {
+    if (!myGlContext->IsGlGreaterEqual (2, 1))
+    {
+      return Standard_False;
+    }
+
+    glBindTexture (GL_TEXTURE_RECTANGLE, *myRaytraceOutputTexture);
+
+    glGetTexLevelParameteriv (GL_TEXTURE_RECTANGLE, 0, GL_TEXTURE_WIDTH,  &aSizeX);
+    glGetTexLevelParameteriv (GL_TEXTURE_RECTANGLE, 0, GL_TEXTURE_HEIGHT, &aSizeY);
+
+    toResize = (aSizeX != theSizeX) || (aSizeY != theSizeY);
+    if (toResize)
+    {
+      glDeleteTextures (2, myRaytraceOutputTexture);
+    }
+  }
+  if (!toResize)
+  {
+    return Standard_True;
+  }
+
+  glGenTextures (2, myRaytraceOutputTexture);
+  for (int aTexIter = 0; aTexIter < 2; ++aTexIter)
+  {
+    glBindTexture (GL_TEXTURE_RECTANGLE, myRaytraceOutputTexture[aTexIter]);
+
+    glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_WRAP_S, GL_CLAMP);
+    glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_WRAP_T, GL_CLAMP);
+    glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_WRAP_R, GL_CLAMP);
+
+    glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
+    glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
+
+    glTexImage2D (GL_TEXTURE_RECTANGLE, 0, GL_RGBA32F,
+                  theSizeX, theSizeY, 0,
+                  GL_RGBA, GL_FLOAT, NULL);
+  }
+
+  cl_int anError = CL_SUCCESS;
+
+  if (myRaytraceOutputImage != NULL)
+  {
+    clReleaseMemObject (myRaytraceOutputImage);
+  }
+  if (myRaytraceOutputImageSmooth != NULL)
+  {
+    clReleaseMemObject (myRaytraceOutputImageSmooth);
+  }
+
+  myRaytraceOutputImage = clCreateFromGLTexture2D (myComputeContext, CL_MEM_READ_WRITE,
+                                                   GL_TEXTURE_RECTANGLE, 0,
+                                                   myRaytraceOutputTexture[0], &anError);
+  if (anError != CL_SUCCESS)
+  {
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "Error! Failed to create output image!" << std::endl;
+#endif
+    return Standard_False;
+  }
+
+  myRaytraceOutputImageSmooth = clCreateFromGLTexture2D (myComputeContext, CL_MEM_READ_WRITE,
+                                                         GL_TEXTURE_RECTANGLE, 0,
+                                                         myRaytraceOutputTexture[1], &anError);
+  if (anError != CL_SUCCESS)
+  {
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "Error! Failed to create anti-aliased output image!" << std::endl;
+#endif
+    return Standard_False;
+  }
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : WriteRaytraceSceneToDevice
+// purpose  : Writes scene geometry to OpenCl device
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice()
+{
+  if (myComputeContext == NULL)
+    return Standard_False;
+
+  cl_int anError = CL_SUCCESS;
+
+  if (myRaytraceNormalBuffer != NULL)
+    anError |= clReleaseMemObject (myRaytraceNormalBuffer);
+
+  if (myRaytraceVertexBuffer != NULL)
+    anError |= clReleaseMemObject (myRaytraceVertexBuffer);
+
+  if (myRaytraceTriangleBuffer != NULL)
+    anError |= clReleaseMemObject (myRaytraceTriangleBuffer);
+
+  if (myRaytraceNodeMinPointBuffer != NULL)
+    anError |= clReleaseMemObject (myRaytraceNodeMinPointBuffer);
+
+  if (myRaytraceNodeMaxPointBuffer != NULL)
+    anError |= clReleaseMemObject (myRaytraceNodeMaxPointBuffer);
+
+  if (myRaytraceNodeDataRcrdBuffer != NULL)
+    anError |= clReleaseMemObject (myRaytraceNodeDataRcrdBuffer);
+
+  if (myRaytraceMaterialBuffer != NULL)
+    anError |= clReleaseMemObject (myRaytraceMaterialBuffer);
+
+  if (anError != CL_SUCCESS)
+  {
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "Error! Failed to release OpenCL scene buffers!" << std::endl;
+#endif
+    return Standard_False;
+  }
+
+  // Create geometry buffers
+  cl_int anErrorTemp = CL_SUCCESS;
+  const size_t myVertexBufferSize = myRaytraceSceneData.Vertices.size() > 0
+                                  ? myRaytraceSceneData.Vertices.size() : 1;
+
+  myRaytraceVertexBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
+                                           myVertexBufferSize * sizeof(cl_float4), NULL, &anErrorTemp);
+  anError |= anErrorTemp;
+
+  const size_t myNormalBufferSize = myRaytraceSceneData.Normals.size() > 0
+                                  ? myRaytraceSceneData.Normals.size() : 1;
+  myRaytraceNormalBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
+                                           myNormalBufferSize * sizeof(cl_float4), NULL, &anErrorTemp);
+  anError |= anErrorTemp;
+
+  const size_t myTriangleBufferSize = myRaytraceSceneData.Triangles.size() > 0
+                                    ? myRaytraceSceneData.Triangles.size() : 1;
+  myRaytraceTriangleBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
+                                             myTriangleBufferSize * sizeof(cl_int4), NULL, &anErrorTemp);
+  anError |= anErrorTemp;
+  if (anError != CL_SUCCESS)
+  {
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "Error! Failed to create OpenCL geometry buffers!" << std::endl;
+#endif
+    return Standard_False;
+  }
+
+  // Create material buffer
+  const size_t myMaterialBufferSize = myRaytraceSceneData.Materials.size() > 0
+                                    ? myRaytraceSceneData.Materials.size() : 1;
+  myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
+                                             myMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL,
+                                             &anErrorTemp);
+  if (anErrorTemp != CL_SUCCESS)
+  {
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "Error! Failed to create OpenCL material buffer!" << std::endl;
+#endif
+    return Standard_False;
+  }
+
+  // Create BVH buffers
+  OpenGl_BVH aTree = myBVHBuilder.Tree();
+  const size_t myNodeMinPointBufferSize = aTree.MinPointBuffer().size() > 0
+                                        ? aTree.MinPointBuffer().size() : 1;
+  myRaytraceNodeMinPointBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
+                                                 myNodeMinPointBufferSize * sizeof(cl_float4), NULL,
+                                                 &anErrorTemp);
+  anError |= anErrorTemp;
+
+  const size_t myNodeMaxPointBufferSize = aTree.MaxPointBuffer().size() > 0
+                                        ? aTree.MaxPointBuffer().size() : 1;
+  myRaytraceNodeMaxPointBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
+                                                 myNodeMaxPointBufferSize * sizeof(cl_float4), NULL,
+                                                 &anError);
+  anError |= anErrorTemp;
+
+  const size_t myNodeDataRecordBufferSize = aTree.DataRcrdBuffer().size() > 0
+                                          ? aTree.DataRcrdBuffer().size() : 1;
+  myRaytraceNodeDataRcrdBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
+                                                 myNodeDataRecordBufferSize * sizeof(cl_int4), NULL,
+                                                 &anError);
+  anError |= anErrorTemp;
+  if (anError != CL_SUCCESS)
+  {
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "Error! Failed to create OpenCL BVH buffers!" << std::endl;
+#endif
+    return Standard_False;
+  }
+
+  // Write scene geometry buffers
+  if (myRaytraceSceneData.Triangles.size() > 0)
+  {
+    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceVertexBuffer, CL_FALSE,
+                                     0, myRaytraceSceneData.Vertices.size() * sizeof(cl_float4),
+                                     &myRaytraceSceneData.Vertices.front(),
+                                     0, NULL, NULL);
+    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNormalBuffer, CL_FALSE,
+                                     0, myRaytraceSceneData.Normals.size() * sizeof(cl_float4),
+                                     &myRaytraceSceneData.Normals.front(),
+                                     0, NULL, NULL);
+    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceTriangleBuffer, CL_FALSE,
+                                     0, myRaytraceSceneData.Triangles.size() * sizeof(cl_int4),
+                                     &myRaytraceSceneData.Triangles.front(),
+                                     0, NULL, NULL);
+    if (anError != CL_SUCCESS)
+    {
+  #ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error! Failed to write OpenCL geometry buffers!" << std::endl;
+  #endif
+      return Standard_False;
+    }
+  }
+
+  // Write BVH buffers
+  if (aTree.DataRcrdBuffer().size() > 0)
+  {
+    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeMinPointBuffer, CL_FALSE,
+                                     0, aTree.MinPointBuffer().size() * sizeof(cl_float4),
+                                     &aTree.MinPointBuffer().front(),
+                                     0, NULL, NULL);
+    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeMaxPointBuffer, CL_FALSE,
+                                     0, aTree.MaxPointBuffer().size() * sizeof(cl_float4),
+                                     &aTree.MaxPointBuffer().front(),
+                                     0, NULL, NULL);
+    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeDataRcrdBuffer, CL_FALSE,
+                                     0, aTree.DataRcrdBuffer().size() * sizeof(cl_int4),
+                                     &aTree.DataRcrdBuffer().front(),
+                                     0, NULL, NULL);
+    if (anError != CL_SUCCESS)
+    {
+  #ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error! Failed to write OpenCL BVH buffers!" << std::endl;
+  #endif
+      return Standard_False;
+    }
+  }
+
+  // Write material buffers
+  if (myRaytraceSceneData.Materials.size() > 0)
+  {
+    const size_t aSize    = myRaytraceSceneData.Materials.size();
+    const void*  aDataPtr = myRaytraceSceneData.Materials.front().Packed();
+
+    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceMaterialBuffer, CL_FALSE,
+                                     0, aSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr,
+                                     0, NULL, NULL);
+    if (anError != CL_SUCCESS)
+    {
+  #ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error! Failed to write OpenCL material buffer!" << std::endl;
+  #endif
+      return Standard_False;
+    }
+  }
+
+  anError |= clFinish (myRaytraceQueue);
+#ifdef RAY_TRACE_PRINT_INFO
+  if (anError != CL_SUCCESS)
+    std::cout << "Error! Failed to set scene data buffers!" << std::endl;
+#endif
+
+  if (anError == CL_SUCCESS)
+    myIsRaytraceDataValid = myRaytraceSceneData.Triangles.size() > 0;
+
+#ifdef RAY_TRACE_PRINT_INFO
+
+  float aMemUsed = static_cast<float> (
+    myRaytraceSceneData.Materials.size()) * sizeof (OpenGl_RaytraceMaterial);
+
+  aMemUsed += static_cast<float> (
+    myRaytraceSceneData.Triangles.size() * sizeof (OpenGl_RTVec4i) +
+    myRaytraceSceneData.Vertices.size()  * sizeof (OpenGl_RTVec4f) +
+    myRaytraceSceneData.Normals.size()   * sizeof (OpenGl_RTVec4f));
+
+  aMemUsed += static_cast<float> (
+    aTree.MinPointBuffer().size() * sizeof (OpenGl_RTVec4f) +
+    aTree.MaxPointBuffer().size() * sizeof (OpenGl_RTVec4f) +
+    aTree.DataRcrdBuffer().size() * sizeof (OpenGl_RTVec4i));
+
+  std::cout << "GPU memory used (Mb): " << aMemUsed / 1e6f << std::endl;
+
+#endif
+
+  myRaytraceSceneData.Clear();
+
+  myBVHBuilder.CleanUp();
+
+  return (CL_SUCCESS == anError);
+}
+
+#define OPENCL_GROUP_SIZE_TEST_
+
+// =======================================================================
+// function : RunRaytraceOpenCLKernels
+// purpose  : Runs OpenCL ray-tracing kernels
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView,
+                                                             const GLfloat          theOrigins[16],
+                                                             const GLfloat          theDirects[16],
+                                                             const int              theSizeX,
+                                                             const int              theSizeY)
+{
+  if (myRaytraceRenderKernel == NULL || myRaytraceQueue == NULL)
+    return Standard_False;
+
+  ////////////////////////////////////////////////////////////
+  // Set kernel arguments
+
+  cl_uint anIndex = 0;
+  cl_int  anError = 0;
+
+  anError  = clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceOutputImage);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceEnvironment);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceNodeMinPointBuffer);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceNodeMaxPointBuffer);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceNodeDataRcrdBuffer);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceLightSourceBuffer);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceMaterialBuffer);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceVertexBuffer);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceNormalBuffer);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_mem), &myRaytraceTriangleBuffer);
+
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_float16), theOrigins);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_float16), theDirects);
+
+  cl_int aLightCount =  static_cast<cl_int> (myRaytraceSceneData.LightSources.size());
+
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_int),   &aLightCount);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_float), &myRaytraceSceneEpsilon);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_float), &myRaytraceSceneRadius);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_int),   &theCView.IsShadowsEnabled);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_int),   &theCView.IsReflectionsEnabled);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_int),   &theSizeX);
+  anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
+                             sizeof(cl_int),   &theSizeY);
+  if (anError != CL_SUCCESS)
+  {
+    const TCollection_ExtendedString aMsg = "Error! Failed to set arguments of ray-tracing kernel!";
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                              GL_DEBUG_TYPE_ERROR_ARB,
+                              0,
+                              GL_DEBUG_SEVERITY_HIGH_ARB,
+                              aMsg);
+    return Standard_False;
+  }
+
+  // Note: second-pass 'smoothing' kernel runs only if anti-aliasing is enabled
+  if (theCView.IsAntialiasingEnabled)
+  {
+    anIndex = 0;
+    anError  = clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceOutputImage);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceOutputImageSmooth);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceEnvironment);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceNodeMinPointBuffer);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceNodeMaxPointBuffer);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceNodeDataRcrdBuffer);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceLightSourceBuffer);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceMaterialBuffer);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceVertexBuffer);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceNormalBuffer);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_mem), &myRaytraceTriangleBuffer);
+
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_float16), theOrigins);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                                sizeof(cl_float16), theDirects);
+
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_int),   &aLightCount);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_float), &myRaytraceSceneEpsilon);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_float), &myRaytraceSceneRadius);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_int),   &theCView.IsShadowsEnabled);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_int),   &theCView.IsReflectionsEnabled);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_int),   &theSizeX);
+    anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
+                               sizeof(cl_int),   &theSizeY);
+    if (anError != CL_SUCCESS)
+    {
+      const TCollection_ExtendedString aMsg = "Error! Failed to set arguments of 'smoothing' kernel!";
+      myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                                GL_DEBUG_TYPE_ERROR_ARB,
+                                0,
+                                GL_DEBUG_SEVERITY_HIGH_ARB,
+                                aMsg);
+      return Standard_False;
+    }
+  }
+
+  // Set work size
+  size_t aLocSizeRender[] = { myIsAmdComputePlatform ? 2 : 4, 32 };
+
+#ifdef OPENCL_GROUP_SIZE_TEST
+  for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1 )
+  for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1 )
+#endif
+  {
+#ifdef OPENCL_GROUP_SIZE_TEST
+    aLocSizeRender[0] = aLocX;
+    aLocSizeRender[1] = aLocY;
+#endif
+
+    size_t aWorkSizeX = theSizeX;
+    if (aWorkSizeX % aLocSizeRender[0] != 0)
+      aWorkSizeX += aLocSizeRender[0] - aWorkSizeX % aLocSizeRender[0];
+
+    size_t aWokrSizeY = theSizeY;
+    if (aWokrSizeY % aLocSizeRender[1] != 0 )
+      aWokrSizeY += aLocSizeRender[1] - aWokrSizeY % aLocSizeRender[1];
+
+    size_t aGlbSizeRender[] = { aWorkSizeX, aWokrSizeY };
+
+    // Run kernel
+    cl_event anEvent (NULL), anEventSmooth (NULL);
+    anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceRenderKernel,
+                                      2, NULL, aGlbSizeRender, aLocSizeRender,
+                                      0, NULL, &anEvent);
+    if (anError != CL_SUCCESS)
+    {
+      const TCollection_ExtendedString aMsg = "Error! Failed to execute the ray-tracing kernel!";
+      myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                                GL_DEBUG_TYPE_ERROR_ARB,
+                                0,
+                                GL_DEBUG_SEVERITY_HIGH_ARB,
+                                aMsg);
+      return Standard_False;
+    }
+    clWaitForEvents (1, &anEvent);
+
+    if (theCView.IsAntialiasingEnabled)
+    {
+      size_t aLocSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4,
+                                  myIsAmdComputePlatform ? 8 : 32 };
+
+#ifdef OPENCL_GROUP_SIZE_TEST
+      aLocSizeSmooth[0] = aLocX;
+      aLocSizeSmooth[1] = aLocY;
+#endif
+
+      aWorkSizeX = theSizeX;
+      if (aWorkSizeX % aLocSizeSmooth[0] != 0)
+        aWorkSizeX += aLocSizeSmooth[0] - aWorkSizeX % aLocSizeSmooth[0];
+
+      size_t aWokrSizeY = theSizeY;
+      if (aWokrSizeY % aLocSizeSmooth[1] != 0 )
+        aWokrSizeY += aLocSizeSmooth[1] - aWokrSizeY % aLocSizeSmooth[1];
+
+      size_t aGlbSizeSmooth [] = { aWorkSizeX, aWokrSizeY };
+      anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceSmoothKernel,
+                                        2, NULL, aGlbSizeSmooth, aLocSizeSmooth,
+                                        0, NULL, &anEventSmooth);
+      clWaitForEvents (1, &anEventSmooth);
+
+      if (anError != CL_SUCCESS)
+      {
+        const TCollection_ExtendedString aMsg = "Error! Failed to execute the 'smoothing' kernel!";
+        myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+                                  GL_DEBUG_TYPE_ERROR_ARB,
+                                  0,
+                                  GL_DEBUG_SEVERITY_HIGH_ARB,
+                                  aMsg);
+        return Standard_False;
+      }
+    }
+
+    // Get the profiling data
+#if defined (RAY_TRACE_PRINT_INFO) || defined(OPENCL_GROUP_SIZE_TEST)
+
+    cl_ulong aTimeStart,
+             aTimeFinal;
+
+    clGetEventProfilingInfo (anEvent, CL_PROFILING_COMMAND_START,
+                             sizeof(aTimeStart), &aTimeStart, NULL);
+    clGetEventProfilingInfo (anEvent, CL_PROFILING_COMMAND_END,
+                             sizeof(aTimeFinal), &aTimeFinal, NULL);
+    std::cout << "\tRender time (ms): " << ( aTimeFinal - aTimeStart ) / 1e6f << std::endl;
+
+    if (theCView.IsAntialiasingEnabled)
+    {
+      clGetEventProfilingInfo (anEventSmooth, CL_PROFILING_COMMAND_START,
+                               sizeof(aTimeStart), &aTimeStart, NULL);
+      clGetEventProfilingInfo (anEventSmooth, CL_PROFILING_COMMAND_END,
+                               sizeof(aTimeFinal), &aTimeFinal, NULL);
+      std::cout << "\tSmoothing time (ms): " << ( aTimeFinal - aTimeStart ) / 1e6f << std::endl;
+    }
+#endif
+
+    if (anEvent != NULL)
+      clReleaseEvent (anEvent);
+
+    if (anEventSmooth != NULL)
+      clReleaseEvent (anEventSmooth);
+  }
+
+  return Standard_True;
+}
+
+// =======================================================================
+// function : ComputeInverseMatrix
+// purpose  : Computes inversion of 4x4 floating-point matrix
+// =======================================================================
+template <typename T>
+void ComputeInverseMatrix (const T m[16], T inv[16])
+{
+  inv[ 0] = m[ 5] * (m[10] * m[15] - m[11] * m[14]) -
+            m[ 9] * (m[ 6] * m[15] - m[ 7] * m[14]) -
+            m[13] * (m[ 7] * m[10] - m[ 6] * m[11]);
+
+  inv[ 1] = m[ 1] * (m[11] * m[14] - m[10] * m[15]) -
+            m[ 9] * (m[ 3] * m[14] - m[ 2] * m[15]) -
+            m[13] * (m[ 2] * m[11] - m[ 3] * m[10]);
+
+  inv[ 2] = m[ 1] * (m[ 6] * m[15] - m[ 7] * m[14]) -
+            m[ 5] * (m[ 2] * m[15] - m[ 3] * m[14]) -
+            m[13] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]);
+
+  inv[ 3] = m[ 1] * (m[ 7] * m[10] - m[ 6] * m[11]) -
+            m[ 5] * (m[ 3] * m[10] - m[ 2] * m[11]) -
+            m[ 9] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]);
+
+  inv[ 4] = m[ 4] * (m[11] * m[14] - m[10] * m[15]) -
+            m[ 8] * (m[ 7] * m[14] - m[ 6] * m[15]) -
+            m[12] * (m[ 6] * m[11] - m[ 7] * m[10]);
+
+  inv[ 5] = m[ 0] * (m[10] * m[15] - m[11] * m[14]) -
+            m[ 8] * (m[ 2] * m[15] - m[ 3] * m[14]) -
+            m[12] * (m[ 3] * m[10] - m[ 2] * m[11]);
+
+  inv[ 6] = m[ 0] * (m[ 7] * m[14] - m[ 6] * m[15]) -
+            m[ 4] * (m[ 3] * m[14] - m[ 2] * m[15]) -
+            m[12] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]);
+
+  inv[ 7] = m[ 0] * (m[ 6] * m[11] - m[ 7] * m[10]) -
+            m[ 4] * (m[ 2] * m[11] - m[ 3] * m[10]) -
+            m[ 8] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]);
+
+  inv[ 8] = m[ 4] * (m[ 9] * m[15] - m[11] * m[13]) -
+            m[ 8] * (m[ 5] * m[15] - m[ 7] * m[13]) -
+            m[12] * (m[ 7] * m[ 9] - m[ 5] * m[11]);
+
+  inv[ 9] = m[ 0] * (m[11] * m[13] - m[ 9] * m[15]) -
+            m[ 8] * (m[ 3] * m[13] - m[ 1] * m[15]) -
+            m[12] * (m[ 1] * m[11] - m[ 3] * m[ 9]);
+
+  inv[10] = m[ 0] * (m[ 5] * m[15] - m[ 7] * m[13]) -
+            m[ 4] * (m[ 1] * m[15] - m[ 3] * m[13]) -
+            m[12] * (m[ 3] * m[ 5] - m[ 1] * m[ 7]);
+
+  inv[11] = m[ 0] * (m[ 7] * m[ 9] - m[ 5] * m[11]) -
+            m[ 4] * (m[ 3] * m[ 9] - m[ 1] * m[11]) -
+            m[ 8] * (m[ 1] * m[ 7] - m[ 3] * m[ 5]);
+
+  inv[12] = m[ 4] * (m[10] * m[13] - m[ 9] * m[14]) -
+            m[ 8] * (m[ 6] * m[13] - m[ 5] * m[14]) -
+            m[12] * (m[ 5] * m[10] - m[ 6] * m[ 9]);
+
+  inv[13] = m[ 0] * (m[ 9] * m[14] - m[10] * m[13]) -
+            m[ 8] * (m[ 1] * m[14] - m[ 2] * m[13]) -
+            m[12] * (m[ 2] * m[ 9] - m[ 1] * m[10]);
+
+  inv[14] = m[ 0] * (m[ 6] * m[13] - m[ 5] * m[14]) -
+            m[ 4] * (m[ 2] * m[13] - m[ 1] * m[14]) -
+            m[12] * (m[ 1] * m[ 6] - m[ 2] * m[ 5]);
+
+  inv[15] = m[ 0] * (m[ 5] * m[10] - m[ 6] * m[ 9]) -
+            m[ 4] * (m[ 1] * m[10] - m[ 2] * m[ 9]) -
+            m[ 8] * (m[ 2] * m[ 5] - m[ 1] * m[ 6]);
+
+  T det = m[0] * inv[ 0] +
+          m[1] * inv[ 4] +
+          m[2] * inv[ 8] +
+          m[3] * inv[12];
+
+  if (det == T (0.0)) return;
+
+  det = T (1.0) / det;
+
+  for (int i = 0; i < 16; ++i)
+    inv[i] *= det;
+}
+
+// =======================================================================
+// function : GenerateCornerRays
+// purpose  : Generates primary rays for corners of screen quad
+// =======================================================================
+void GenerateCornerRays (const GLdouble theInvModelProj[16],
+                         float          theOrigins[16],
+                         float          theDirects[16])
+{
+  int aOriginIndex = 0;
+  int aDirectIndex = 0;
+
+  for (int y = -1; y <= 1; y += 2)
+  {
+    for (int x = -1; x <= 1; x += 2)
+    {
+      OpenGl_RTVec4f aOrigin (float(x),
+                              float(y),
+                              -1.f,
+                              1.f);
+
+      aOrigin = MatVecMult (theInvModelProj, aOrigin);
+
+      OpenGl_RTVec4f aDirect (float(x),
+                              float(y),
+                              1.f,
+                              1.f);
+
+      aDirect = MatVecMult (theInvModelProj, aDirect) - aOrigin;
+
+      GLdouble aInvLen = 1.f / sqrt (aDirect.x() * aDirect.x() +
+                                     aDirect.y() * aDirect.y() +
+                                     aDirect.z() * aDirect.z());
+
+      theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.x());
+      theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.y());
+      theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.z());
+      theOrigins [aOriginIndex++] = 1.f;
+
+      theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.x() * aInvLen);
+      theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.y() * aInvLen);
+      theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.z() * aInvLen);
+      theDirects [aDirectIndex++] = 0.f;
+    }
+  }
+}
+
+// =======================================================================
+// function : Raytrace
+// purpose  : Redraws the window using OpenCL ray tracing
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView,
+                                             const int              theSizeX,
+                                             int                    theSizeY,
+                                             const Tint             theToSwap)
+{
+  if (!InitOpenCL())
+    return Standard_False;
+
+  if (!ResizeRaytraceOutputBuffer (theSizeX, theSizeY))
+    return Standard_False;
+
+  if (!UpdateRaytraceEnvironmentMap())
+    return Standard_False;
+
+  if (!UpdateRaytraceGeometry (Standard_True))
+    return Standard_False;
+
+  // Get model-view and projection matrices
+  TColStd_Array2OfReal theOrientation (0, 3, 0, 3);
+  TColStd_Array2OfReal theViewMapping (0, 3, 0, 3);
+
+  myView->GetMatrices (theOrientation, theViewMapping, Standard_True);
+
+  GLdouble aOrientationMatrix[16];
+  GLdouble aViewMappingMatrix[16];
+  GLdouble aOrientationInvers[16];
+
+  for (int j = 0; j < 4; ++j)
+    for (int i = 0; i < 4; ++i)
+    {
+      aOrientationMatrix [4 * j + i] = theOrientation (i, j);
+      aViewMappingMatrix [4 * j + i] = theViewMapping (i, j);
+    }
+
+  ComputeInverseMatrix (aOrientationMatrix, aOrientationInvers);
+
+  if (!UpdateRaytraceLightSources (aOrientationInvers))
+    return Standard_False;
+
+  // Generate primary rays for corners of the screen quad
+  glMatrixMode (GL_MODELVIEW);
+
+  glLoadMatrixd (aViewMappingMatrix);
+  glMultMatrixd (aOrientationMatrix);
+
+  GLdouble aModelProject[16];
+  GLdouble aInvModelProj[16];
+
+  glGetDoublev (GL_MODELVIEW_MATRIX, aModelProject);
+
+  ComputeInverseMatrix (aModelProject, aInvModelProj);
+
+  GLfloat aOrigins[16];
+  GLfloat aDirects[16];
+
+  GenerateCornerRays (aInvModelProj,
+                      aOrigins,
+                      aDirects);
+
+  // Compute ray-traced image using OpenCL kernel
+  cl_mem anImages[] = { myRaytraceOutputImage, myRaytraceOutputImageSmooth };
+  cl_int anError = clEnqueueAcquireGLObjects (myRaytraceQueue,
+                                              2, anImages,
+                                              0, NULL, NULL);
+  clFinish (myRaytraceQueue);
+
+  if (myIsRaytraceDataValid)
+  {
+    RunRaytraceOpenCLKernels (theCView,
+                              aOrigins,
+                              aDirects,
+                              theSizeX,
+                              theSizeY);
+  }
+
+  anError |= clEnqueueReleaseGLObjects (myRaytraceQueue,
+                                        2, anImages,
+                                        0, NULL, NULL);
+  clFinish (myRaytraceQueue);
+
+  // Draw background
+  glPushAttrib (GL_ENABLE_BIT |
+                GL_CURRENT_BIT |
+                GL_COLOR_BUFFER_BIT |
+                GL_DEPTH_BUFFER_BIT);
+
+  glDisable (GL_DEPTH_TEST);
+
+  if (NamedStatus & OPENGL_NS_WHITEBACK)
+  {
+    glClearColor (1.0f, 1.0f, 1.0f, 1.0f);
+  }
+  else
+  {
+    glClearColor (myBgColor.rgb[0],
+                  myBgColor.rgb[1],
+                  myBgColor.rgb[2],
+                  1.0f);
+  }
+
+  glClear (GL_COLOR_BUFFER_BIT);
+
+  Handle(OpenGl_Workspace) aWorkspace (this);
+  myView->DrawBackground (aWorkspace);
+
+  // Draw dummy quad to show result image
+  glEnable (GL_COLOR_MATERIAL);
+  glEnable (GL_BLEND);
+
+  glDisable (GL_DEPTH_TEST);
+
+  glBlendFunc (GL_ONE, GL_SRC_ALPHA);
+
+  glEnable (GL_TEXTURE_RECTANGLE);
+
+  glMatrixMode (GL_PROJECTION);
+  glLoadIdentity();
+
+  glMatrixMode (GL_MODELVIEW);
+  glLoadIdentity();
+
+  glColor3f (1.0f, 1.0f, 1.0f);
+
+  glBindTexture (GL_TEXTURE_RECTANGLE, myRaytraceOutputTexture[theCView.IsAntialiasingEnabled ? 1 : 0]);
+
+  if (myIsRaytraceDataValid)
+  {
+    glBegin (GL_QUADS);
+    {
+      glTexCoord2i (       0,        0);   glVertex2f (-1.f, -1.f);
+      glTexCoord2i (       0, theSizeY);   glVertex2f (-1.f,  1.f);
+      glTexCoord2i (theSizeX, theSizeY);   glVertex2f ( 1.f,  1.f);
+      glTexCoord2i (theSizeX,        0);   glVertex2f ( 1.f, -1.f);
+    }
+    glEnd();
+  }
+
+  glPopAttrib();
+
+  // Swap the buffers
+  if (theToSwap)
+  {
+    GetGlContext()->SwapBuffers();
+    myBackBufferRestored = Standard_False;
+  }
+  else
+    glFlush();
+
+  return Standard_True;
+}
+
+#endif
index 0ead3de..058f572 100755 (executable)
@@ -5,6 +5,7 @@ CSF_OpenGlLibs
 CSF_objc
 CSF_Appkit
 CSF_IOKit
+CSF_OPENCL
 CSF_FREETYPE
 CSF_GL2PS
 CSF_user32
index e9543ea..a21d806 100755 (executable)
@@ -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 ;
index 06212e9..5498d2e 100755 (executable)
@@ -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;
index c885327..e32241f 100755 (executable)
@@ -19,7 +19,7 @@
 // and conditions governing the rights and limitations under the License.
 
 #ifdef HAVE_CONFIG_H
-# include <config.h>
+  #include <config.h>
 #endif
 
 #include <OpenGl_GlCore20.hxx>
@@ -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);
 }
index eab3eb5..466a2c0 100755 (executable)
@@ -1,4 +1,5 @@
 FAILED /\bFaulty\b/ bad shape
 IGNORE /^Error [23]d = [\d.-]+/ debug output of blend command
 SKIPPED /Error: unsupported locale specification/ locale is unavailable on tested system
+SKIPPED /OCCT was compiled without OpenCL support!/
 OK /Relative error of mass computation/ message from vprops
diff --git a/tests/bugs/vis/bug24130 b/tests/bugs/vis/bug24130
new file mode 100644 (file)
index 0000000..7c30ba3
--- /dev/null
@@ -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
+  }
+}