0024795: TKOpenGl - new compilation warnings in OpenGl_Workspace_Raytrace.cxx
[occt.git] / src / OpenGl / OpenGl_Workspace_Raytrace.cxx
old mode 100644 (file)
new mode 100755 (executable)
index 192b21e..b03caad
@@ -2,47 +2,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.
+// This file is part of Open CASCADE Technology software library.
 //
-// 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.
+// This library is free software; you can redistribute it and/or modify it under
+// the terms of the GNU Lesser General Public License version 2.1 as published
+// by the Free Software Foundation, with special exception defined in the file
+// OCCT_LGPL_EXCEPTION.txt. Consult the file LICENSE_LGPL_21.txt included in OCCT
+// distribution for complete text of the license and disclaimer of any warranty.
 //
-// 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 <OpenGl_Cl.hxx>
-
-#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
+// Alternatively, this file may be used under the terms of Open CASCADE
+// commercial license or contractual agreement.
 
-#include <OpenGl_Context.hxx>
+#include <NCollection_Mat4.hxx>
+#include <OpenGl_ArbFBO.hxx>
+#include <OpenGl_FrameBuffer.hxx>
 #include <OpenGl_Texture.hxx>
+#include <OpenGl_VertexBuffer.hxx>
 #include <OpenGl_View.hxx>
 #include <OpenGl_Workspace.hxx>
+#include <OSD_File.hxx>
+#include <OSD_Protection.hxx>
+#include <Standard_Assert.hxx>
 
 using namespace OpenGl_Raytrace;
 
@@ -53,17 +33,14 @@ using namespace OpenGl_Raytrace;
   #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)
+template<typename T>
+BVH_Vec4f MatVecMult (const T m[16], const BVH_Vec4f& v)
 {
-  return OpenGl_RTVec4f (
+  return BVH_Vec4f (
     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() +
@@ -74,87 +51,6 @@ OpenGl_RTVec4f MatVecMult (const T m[16], const OpenGl_RTVec4f& v)
                         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
@@ -169,7 +65,7 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theC
 
   if (!theCheck)
   {
-    myRaytraceSceneData.Clear();
+    myRaytraceGeometry.Clear();
 
     myIsRaytraceDataValid = Standard_False;
   }
@@ -181,7 +77,7 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theC
     }
   }
 
-  float* aTransform (NULL);
+  Standard_ShortReal* aTransform (NULL);
 
   // The set of processed structures (reflected to ray-tracing)
   // This set is used to remove out-of-date records from the
@@ -199,7 +95,7 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theC
 
     const OpenGl_ArrayOfStructure& aStructArray = aPriorityList.ArrayOfStructures();
 
-    for (int anIndex = 0; anIndex < aStructArray.Length(); ++anIndex)
+    for (Standard_Integer anIndex = 0; anIndex < aStructArray.Length(); ++anIndex)
     {
       OpenGl_SequenceOfStructure::Iterator aStructIt;
 
@@ -222,10 +118,10 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theC
           if (aStructure->Transformation()->mat != NULL)
           {
             if (aTransform == NULL)
-              aTransform = new float[16];
+              aTransform = new Standard_ShortReal[16];
 
-            for (int i = 0; i < 4; ++i)
-              for (int j = 0; j < 4; ++j)
+            for (Standard_Integer i = 0; i < 4; ++i)
+              for (Standard_Integer j = 0; j < 4; ++j)
               {
                 aTransform[j * 4 + i] = aStructure->Transformation()->mat[i][j];
               }
@@ -257,32 +153,22 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theC
     // Actualize OpenGL layer list state
     myLayersModificationStatus = myView->LayerList().ModificationState();
 
+    // Rebuild bottom-level and high-level BVHs
+    myRaytraceGeometry.ProcessAcceleration();
 
-#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 Standard_ShortReal aMinRadius = Max (fabs (myRaytraceGeometry.Box().CornerMin().x()), Max (
+      fabs (myRaytraceGeometry.Box().CornerMin().y()), fabs (myRaytraceGeometry.Box().CornerMin().z())));
+    const Standard_ShortReal aMaxRadius = Max (fabs (myRaytraceGeometry.Box().CornerMax().x()), Max (
+      fabs (myRaytraceGeometry.Box().CornerMax().y()), fabs (myRaytraceGeometry.Box().CornerMax().z())));
 
-    const float aScaleFactor = 1.5f;
+    myRaytraceSceneRadius = 2.f /* scale factor */ * Max (aMinRadius, aMaxRadius);
 
-    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()))) );
+    const BVH_Vec4f aSize = myRaytraceGeometry.Box().Size();
 
-    myRaytraceSceneEpsilon = Max (1e-4f, myRaytraceSceneRadius * 1e-4f);
+    myRaytraceSceneEpsilon = Max (1e-4f, 1e-4f * sqrtf (
+      aSize.x() * aSize.x() + aSize.y() * aSize.y() + aSize.z() * aSize.z()));
 
-    return WriteRaytraceSceneToDevice();
+    return UploadRaytraceData();
   }
 
   delete [] aTransform;
@@ -292,7 +178,7 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theC
 
 // =======================================================================
 // function : CheckRaytraceStructure
-// purpose  : Adds OpenGL structure to ray-traced scene geometry
+// purpose  :  Checks to see if the structure is modified
 // =======================================================================
 Standard_Boolean OpenGl_Workspace::CheckRaytraceStructure (const OpenGl_Structure* theStructure)
 {
@@ -321,40 +207,39 @@ Standard_Boolean OpenGl_Workspace::CheckRaytraceStructure (const OpenGl_Structur
 // function : CreateMaterial
 // purpose  : Creates ray-tracing material properties
 // =======================================================================
-void CreateMaterial (const OPENGL_SURF_PROP&  theProp,
-                     OpenGl_RaytraceMaterial& theMaterial)
+void CreateMaterial (const OPENGL_SURF_PROP& theProp, OpenGl_RaytraceMaterial& theMaterial)
 {
   const float* aSrcAmb = theProp.isphysic ? theProp.ambcol.rgb : theProp.matcol.rgb;
-  theMaterial.Ambient = OpenGl_RTVec4f (aSrcAmb[0] * theProp.amb,
-                                        aSrcAmb[1] * theProp.amb,
-                                        aSrcAmb[2] * theProp.amb,
-                                        1.0f);
+  theMaterial.Ambient = BVH_Vec4f (aSrcAmb[0] * theProp.amb,
+                                   aSrcAmb[1] * theProp.amb,
+                                   aSrcAmb[2] * theProp.amb,
+                                   1.0f);
 
   const float* aSrcDif = theProp.isphysic ? theProp.difcol.rgb : theProp.matcol.rgb;
-  theMaterial.Diffuse = OpenGl_RTVec4f (aSrcDif[0] * theProp.diff,
-                                        aSrcDif[1] * theProp.diff,
-                                        aSrcDif[2] * theProp.diff,
-                                        1.0f);
+  theMaterial.Diffuse = BVH_Vec4f (aSrcDif[0] * theProp.diff,
+                                   aSrcDif[1] * theProp.diff,
+                                   aSrcDif[2] * theProp.diff,
+                                   1.0f);
 
   const float aDefSpecCol[4] = {1.0f, 1.0f, 1.0f, 1.0f};
   const float* aSrcSpe = theProp.isphysic ? theProp.speccol.rgb : aDefSpecCol;
-  theMaterial.Specular = OpenGl_RTVec4f (aSrcSpe[0] * theProp.spec,
-                                         aSrcSpe[1] * theProp.spec,
-                                         aSrcSpe[2] * theProp.spec,
-                                         theProp.shine);
+  theMaterial.Specular = BVH_Vec4f (aSrcSpe[0] * theProp.spec,
+                                    aSrcSpe[1] * theProp.spec,
+                                    aSrcSpe[2] * theProp.spec,
+                                    theProp.shine);
 
   const float* aSrcEms = theProp.isphysic ? theProp.emscol.rgb : theProp.matcol.rgb;
-  theMaterial.Emission = OpenGl_RTVec4f (aSrcEms[0] * theProp.emsv,
-                                         aSrcEms[1] * theProp.emsv,
-                                         aSrcEms[2] * theProp.emsv,
-                                         1.0f);
+  theMaterial.Emission = BVH_Vec4f (aSrcEms[0] * theProp.emsv,
+                                    aSrcEms[1] * theProp.emsv,
+                                    aSrcEms[2] * theProp.emsv,
+                                    1.0f);
 
   // 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);
+  theMaterial.Transparency = BVH_Vec4f (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(),
@@ -362,24 +247,19 @@ void CreateMaterial (const OPENGL_SURF_PROP&  theProp,
 
   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;
+  theMaterial.Reflection = BVH_Vec4f (theProp.speccol.rgb[0] * theProp.spec * aReflectionScale,
+                                      theProp.speccol.rgb[1] * theProp.spec * aReflectionScale,
+                                      theProp.speccol.rgb[2] * theProp.spec * aReflectionScale,
+                                      0.f);
 }
 
 // =======================================================================
 // 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)
+Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure* theStructure,
+  const Standard_ShortReal* 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())
@@ -389,88 +269,80 @@ Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure*
   }
 
   // Get structure material
-  int aStructMatID = -1;
+  Standard_Integer aStructMatID = -1;
 
   if (theStructure->AspectFace() != NULL)
   {
-    aStructMatID = static_cast<int> (myRaytraceSceneData.Materials.size());
+    aStructMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
 
     OpenGl_RaytraceMaterial aStructMaterial;
     CreateMaterial (theStructure->AspectFace()->IntFront(), aStructMaterial);
 
-    myRaytraceSceneData.Materials.push_back (aStructMaterial);
+    myRaytraceGeometry.Materials.push_back (aStructMaterial);
   }
 
-  OpenGl_ListOfGroup::Iterator anItg (theStructure->Groups());
-
-  while (anItg.More())
+  for (OpenGl_Structure::GroupIterator aGroupIter (theStructure->DrawGroups()); aGroupIter.More(); aGroupIter.Next())
   {
     // Get group material
-    int aGroupMatID = -1;
-
-    if (anItg.Value()->AspectFace() != NULL)
+    Standard_Integer aGroupMatID = -1;
+    if (aGroupIter.Value()->AspectFace() != NULL)
     {
-      aGroupMatID = static_cast<int> (myRaytraceSceneData.Materials.size());
+      aGroupMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
 
       OpenGl_RaytraceMaterial aGroupMaterial;
-      CreateMaterial (anItg.Value()->AspectFace()->IntFront(), aGroupMaterial);
+      CreateMaterial (aGroupIter.Value()->AspectFace()->IntFront(), aGroupMaterial);
 
-      myRaytraceSceneData.Materials.push_back (aGroupMaterial);
+      myRaytraceGeometry.Materials.push_back (aGroupMaterial);
     }
 
-    int aMatID = aGroupMatID < 0 ? aStructMatID : aGroupMatID;
+    Standard_Integer aMatID = aGroupMatID < 0 ? aStructMatID : aGroupMatID;
 
-    if (aStructMatID < 0 && aGroupMatID < 0)
+    if (aMatID < 0)
     {
-      aMatID = static_cast<int> (myRaytraceSceneData.Materials.size());
+      aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
 
-      myRaytraceSceneData.Materials.push_back (OpenGl_RaytraceMaterial());
+      myRaytraceGeometry.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)
+    // Add OpenGL elements from group (extract primitives arrays and aspects)
+    for (const OpenGl_ElementNode* aNode = aGroupIter.Value()->FirstNode(); aNode != NULL; aNode = aNode->next)
     {
-      if (TelNil == aNode->type)
+      OpenGl_AspectFace* anAspect = dynamic_cast<OpenGl_AspectFace*> (aNode->elem);
+      if (anAspect != NULL)
       {
-        OpenGl_AspectFace* anAspect = dynamic_cast<OpenGl_AspectFace*> (aNode->elem);
-
-        if (anAspect != NULL)
-        {
-          aMatID = static_cast<int> (myRaytraceSceneData.Materials.size());
+        aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
 
-          OpenGl_RaytraceMaterial aMaterial;
-          CreateMaterial (anAspect->IntFront(), aMaterial);
+        OpenGl_RaytraceMaterial aMaterial;
+        CreateMaterial (anAspect->IntFront(), aMaterial);
 
-          myRaytraceSceneData.Materials.push_back (aMaterial);
-        }
+        myRaytraceGeometry.Materials.push_back (aMaterial);
       }
-      else if (TelParray == aNode->type)
+      else
       {
         OpenGl_PrimitiveArray* aPrimArray = dynamic_cast<OpenGl_PrimitiveArray*> (aNode->elem);
-
         if (aPrimArray != NULL)
         {
-          AddRaytracePrimitiveArray (aPrimArray->PArray(), aMatID, theTransform);
+          NCollection_Handle<BVH_Object<Standard_ShortReal, 4> > aSet =
+            AddRaytracePrimitiveArray (aPrimArray->PArray(), aMatID, theTransform);
+
+          if (!aSet.IsNull())
+            myRaytraceGeometry.Objects().Append (aSet);
         }
       }
     }
-
-    anItg.Next();
   }
 
-  float* aTransform (NULL);
+  Standard_ShortReal* aTransform (NULL);
 
   // Process all connected OpenGL structures
-  OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures());
-
-  while (anIts.More())
+  for (OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures()); anIts.More(); anIts.Next())
   {
     if (anIts.Value()->Transformation()->mat != NULL)
     {
-      float* aTransform = new float[16];
+      Standard_ShortReal* aTransform = new Standard_ShortReal[16];
 
-      for (int i = 0; i < 4; ++i)
-        for (int j = 0; j < 4; ++j)
+      for (Standard_Integer i = 0; i < 4; ++i)
+        for (Standard_Integer j = 0; j < 4; ++j)
         {
           aTransform[j * 4 + i] =
             anIts.Value()->Transformation()->mat[i][j];
@@ -479,8 +351,6 @@ Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure*
 
     if (anIts.Value()->IsRaytracable())
       AddRaytraceStructure (anIts.Value(), aTransform != NULL ? aTransform : theTransform, theElements);
-
-    anIts.Next();
   }
 
   delete[] aTransform;
@@ -494,9 +364,8 @@ Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure*
 // 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)
+OpenGl_TriangleSet* OpenGl_Workspace::AddRaytracePrimitiveArray (
+  const CALL_DEF_PARRAY* theArray, Standard_Integer theMatID, const Standard_ShortReal* theTransform)
 {
   if (theArray->type != TelPolygonsArrayType &&
       theArray->type != TelTrianglesArrayType &&
@@ -505,135 +374,144 @@ Standard_Boolean OpenGl_Workspace::AddRaytracePrimitiveArray (const CALL_DEF_PAR
       theArray->type != TelTriangleStripsArrayType &&
       theArray->type != TelQuadrangleStripsArrayType)
   {
-    return Standard_True;
+    return NULL;
   }
 
   if (theArray->vertices == NULL)
-    return Standard_False;
+    return NULL;
 
 #ifdef RAY_TRACE_PRINT_INFO
   switch (theArray->type)
   {
     case TelPolygonsArrayType:
-      std::cout << "\tTelPolygonsArrayType" << std::endl; break;
+      std::cout << "\tAdding TelPolygonsArrayType" << std::endl; break;
     case TelTrianglesArrayType:
-      std::cout << "\tTelTrianglesArrayType" << std::endl; break;
+      std::cout << "\tAdding TelTrianglesArrayType" << std::endl; break;
     case TelQuadranglesArrayType:
-      std::cout << "\tTelQuadranglesArrayType" << std::endl; break;
+      std::cout << "\tAdding TelQuadranglesArrayType" << std::endl; break;
     case TelTriangleFansArrayType:
-      std::cout << "\tTelTriangleFansArrayType" << std::endl; break;
+      std::cout << "\tAdding TelTriangleFansArrayType" << std::endl; break;
     case TelTriangleStripsArrayType:
-      std::cout << "\tTelTriangleStripsArrayType" << std::endl; break;
+      std::cout << "\tAdding TelTriangleStripsArrayType" << std::endl; break;
     case TelQuadrangleStripsArrayType:
-      std::cout << "\tTelQuadrangleStripsArrayType" << std::endl; break;
+      std::cout << "\tAdding TelQuadrangleStripsArrayType" << 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());
+  OpenGl_TriangleSet* aSet = new OpenGl_TriangleSet;
 
-  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);
+    aSet->Vertices.reserve (theArray->num_vertexs);
 
-    if (theTransform)
-      aVertex = MatVecMult (theTransform, aVertex);
+    for (Standard_Integer aVert = 0; aVert < theArray->num_vertexs; ++aVert)
+    {
+      BVH_Vec4f 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);
+      aSet->Vertices.push_back (aVertex);
+    }
 
-    myRaytraceSceneData.AABB.Add (aVertex);
-  }
+    aSet->Normals.reserve (theArray->num_vertexs);
 
-  myRaytraceSceneData.Normals.reserve (
-    myRaytraceSceneData.Normals.size() + theArray->num_vertexs);
+    for (Standard_Integer aNorm = 0; aNorm < theArray->num_vertexs; ++aNorm)
+    {
+      BVH_Vec4f aNormal;
 
-  for (int aNorm = 0; aNorm < theArray->num_vertexs; ++aNorm)
-  {
-    OpenGl_RTVec4f aNormal;
+      // Note: In case of absence of normals, the
+      // renderer uses generated geometric normals
 
-    // Note: In case of absence of normals, the visualizer
-    // will use generated geometric normals
+      if (theArray->vnormals != NULL)
+      {
+        aNormal = BVH_Vec4f (theArray->vnormals[aNorm].xyz[0],
+                             theArray->vnormals[aNorm].xyz[1],
+                             theArray->vnormals[aNorm].xyz[2],
+                             0.f);
 
-    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);
+      }
 
-      if (theTransform)
-        aNormal = MatVecMult (theTransform, aNormal);
+      aSet->Normals.push_back (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
 
-  if (theArray->num_bounds > 0)
-  {
-#ifdef RAY_TRACE_PRINT_INFO
-    std::cout << "\tNumber of bounds = " << theArray->num_bounds << std::endl;
-#endif
+      Standard_Integer aBoundStart = 0;
+
+      for (Standard_Integer aBound = 0; aBound < theArray->num_bounds; ++aBound)
+      {
+        const Standard_Integer aVertNum = theArray->bounds[aBound];
+
+  #ifdef RAY_TRACE_PRINT_INFO
+        std::cout << "\tAdding indices from bound " << aBound << ": " <<
+                                      aBoundStart << " .. " << aVertNum << std::endl;
+  #endif
 
-    int aVertOffset = 0;
+        if (!AddRaytraceVertexIndices (aSet, theArray, aBoundStart, aVertNum, theMatID))
+        {
+          delete aSet;
+          return NULL;
+        }
 
-    for (int aBound = 0; aBound < theArray->num_bounds; ++aBound)
+        aBoundStart += aVertNum;
+      }
+    }
+    else
     {
-      const int aVertNum = theArray->bounds[aBound];
+      const Standard_Integer aVertNum = theArray->num_edges > 0 ? theArray->num_edges : theArray->num_vertexs;
 
-#ifdef RAY_TRACE_PRINT_INFO
-      std::cout << "\tAdd indices from bound " << aBound << ": " <<
-                                    aVertOffset << ", " << aVertNum << std::endl;
-#endif
+  #ifdef RAY_TRACE_PRINT_INFO
+        std::cout << "\tAdding indices from array: " << aVertNum << std::endl;
+  #endif
 
-      if (!AddRaytraceVertexIndices (theArray, aFirstVert, aVertOffset, aVertNum, theMatID))
+      if (!AddRaytraceVertexIndices (aSet, theArray, 0, aVertNum, theMatID))
       {
-        return Standard_False;
+        delete aSet;
+        return NULL;
       }
-
-      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);
-  }
+  if (aSet->Size() != 0)
+    aSet->MarkDirty();
 
-  return Standard_True;
+  return aSet;
 }
 
 // =======================================================================
 // 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)
+Standard_Boolean OpenGl_Workspace::AddRaytraceVertexIndices (OpenGl_TriangleSet* theSet,
+  const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer 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);
-    default:                           return Standard_False;
+    case TelTrianglesArrayType:
+      return AddRaytraceTriangleArray (theSet, theArray, theOffset, theCount, theMatID);
+
+    case TelQuadranglesArrayType:
+      return AddRaytraceQuadrangleArray (theSet, theArray, theOffset, theCount, theMatID);
+
+    case TelTriangleFansArrayType:
+      return AddRaytraceTriangleFanArray (theSet, theArray, theOffset, theCount, theMatID);
+
+    case TelTriangleStripsArrayType:
+      return AddRaytraceTriangleStripArray (theSet, theArray, theOffset, theCount, theMatID);
+
+    case TelQuadrangleStripsArrayType:
+      return AddRaytraceQuadrangleStripArray (theSet, theArray, theOffset, theCount, theMatID);
+
+    default:
+      return AddRaytracePolygonArray (theSet, theArray, theOffset, theCount, theMatID);
   }
 }
 
@@ -641,33 +519,32 @@ Standard_Boolean OpenGl_Workspace::AddRaytraceVertexIndices (const CALL_DEF_PARR
 // 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)
+Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleArray (OpenGl_TriangleSet* theSet,
+  const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 {
-  if (theVertNum < 3)
+  if (theCount < 3)
     return Standard_True;
 
+  theSet->Elements.reserve (theSet->Elements.size() + theCount / 3);
+
   if (theArray->num_edges > 0)
   {
-    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; aVert += 3)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 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));
+      theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
+                                             theArray->edges[aVert + 1],
+                                             theArray->edges[aVert + 2],
+                                             theMatID));
     }
   }
   else
   {
-    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; aVert += 3)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; aVert += 3)
     {
-      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0,
-                                                               theFirstVert + aVert + 1,
-                                                               theFirstVert + aVert + 2,
-                                                               theMatID));
+      theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
+                                             aVert + 1,
+                                             aVert + 2,
+                                             theMatID));
     }
   }
 
@@ -678,33 +555,32 @@ Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleArray (const CALL_DEF_PARR
 // 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)
+Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleFanArray (OpenGl_TriangleSet* theSet,
+  const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 {
-  if (theVertNum < 3)
+  if (theCount < 3)
     return Standard_True;
 
+  theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
+
   if (theArray->num_edges > 0)
   {
-    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
     {
-      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[theVertOffset],
-                                                               theFirstVert + theArray->edges[aVert + 1],
-                                                               theFirstVert + theArray->edges[aVert + 2],
-                                                               theMatID));
+      theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset],
+                                             theArray->edges[aVert + 1],
+                                             theArray->edges[aVert + 2],
+                                             theMatID));
     }
   }
   else
   {
-    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
     {
-      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset,
-                                                               theFirstVert + aVert + 1,
-                                                               theFirstVert + aVert + 2,
-                                                               theMatID));
+      theSet->Elements.push_back (BVH_Vec4i (theOffset,
+                                             aVert + 1,
+                                             aVert + 2,
+                                             theMatID));
     }
   }
 
@@ -715,45 +591,32 @@ Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleFanArray (const CALL_DEF_P
 // 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)
+Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleStripArray (OpenGl_TriangleSet* theSet,
+  const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 {
-  if (theVertNum < 3)
+  if (theCount < 3)
     return Standard_True;
 
+  theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
+
   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)
+    for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2)
     {
-      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));
+      theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + aCW ? 1 : 0],
+                                             theArray->edges[aVert + aCW ? 0 : 1],
+                                             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)
+    for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2)
     {
-      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + ( aTriNum % 2 ) ? 1 : 0,
-                                                               theFirstVert + aVert + ( aTriNum % 2 ) ? 0 : 1,
-                                                               theFirstVert + aVert + 2,
-                                                               theMatID));
+      theSet->Elements.push_back (BVH_Vec4i (aVert + aCW ? 1 : 0,
+                                             aVert + aCW ? 0 : 1,
+                                             aVert + 2,
+                                             theMatID));
     }
   }
 
@@ -764,43 +627,42 @@ Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleStripArray (const CALL_DEF
 // 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)
+Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleArray (OpenGl_TriangleSet* theSet,
+  const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 {
-  if (theVertNum < 4)
+  if (theCount < 4)
     return Standard_True;
 
+  theSet->Elements.reserve (theSet->Elements.size() + theCount / 2);
+
   if (theArray->num_edges > 0)
   {
-    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 3; aVert += 4)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 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));
+      theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
+                                             theArray->edges[aVert + 1],
+                                             theArray->edges[aVert + 2],
+                                             theMatID));
+
+      theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
+                                             theArray->edges[aVert + 2],
+                                             theArray->edges[aVert + 3],
+                                             theMatID));
     }
   }
   else
   {
-    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 3; aVert += 4)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 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));
+      theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
+                                             aVert + 1,
+                                             aVert + 2,
+                                             theMatID));
+
+      theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
+                                             aVert + 2,
+                                             aVert + 3,
+                                             theMatID));
     }
   }
 
@@ -811,67 +673,42 @@ Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleArray (const CALL_DEF_PA
 // 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)
+Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleStripArray (OpenGl_TriangleSet* theSet,
+  const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 {
-  if (theVertNum < 4)
+  if (theCount < 4)
     return Standard_True;
 
+  theSet->Elements.reserve (theSet->Elements.size() + 2 * theCount - 6);
+
   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)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 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));
+      theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
+                                             theArray->edges[aVert + 1],
+                                             theArray->edges[aVert + 2],
+                                             theMatID));
+
+      theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 1],
+                                             theArray->edges[aVert + 3],
+                                             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)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 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));
+      theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
+                                             aVert + 1,
+                                             aVert + 2,
+                                             theMatID));
+
+      theSet->Elements.push_back (BVH_Vec4i (aVert + 1,
+                                             aVert + 3,
+                                             aVert + 2,
+                                             theMatID));
     }
   }
 
@@ -882,33 +719,32 @@ Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleStripArray (const CALL_D
 // 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)
+Standard_Boolean OpenGl_Workspace::AddRaytracePolygonArray (OpenGl_TriangleSet* theSet,
+  const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 {
-  if (theArray->num_vertexs < 3)
+  if (theCount < 3)
     return Standard_True;
 
-  if (theArray->edges != NULL)
+  theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
+
+  if (theArray->num_edges > 0)
   {
-    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
     {
-      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[theVertOffset],
-                                                               theFirstVert + theArray->edges[aVert + 1],
-                                                               theFirstVert + theArray->edges[aVert + 2],
-                                                               theMatID));
+      theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset],
+                                             theArray->edges[aVert + 1],
+                                             theArray->edges[aVert + 2],
+                                             theMatID));
     }
   }
   else
   {
-    for (int aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+    for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
     {
-      myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset,
-                                                               theFirstVert + aVert + 1,
-                                                               theFirstVert + aVert + 2,
-                                                               theMatID));
+      theSet->Elements.push_back (BVH_Vec4i (theOffset,
+                                             aVert + 1,
+                                             aVert + 2,
+                                             theMatID));
     }
   }
 
@@ -921,1163 +757,1226 @@ Standard_Boolean OpenGl_Workspace::AddRaytracePolygonArray (const CALL_DEF_PARRA
 // =======================================================================
 Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble theInvModelView[16])
 {
-  myRaytraceSceneData.LightSources.clear();
+  myRaytraceGeometry.Sources.clear();
 
-  OpenGl_RTVec4f anAmbient (0.0f, 0.0f, 0.0f, 0.0f);
-  for (OpenGl_ListOfLight::Iterator anItl (myView->LightList());
-       anItl.More(); anItl.Next())
+  myRaytraceGeometry.Ambient = BVH_Vec4f (0.0f, 0.0f, 0.0f, 0.0f);
+
+  for (OpenGl_ListOfLight::Iterator anItl (myView->LightList()); anItl.More(); anItl.Next())
   {
     const OpenGl_Light& aLight = anItl.Value();
+
     if (aLight.Type == Visual3d_TOLS_AMBIENT)
     {
-      anAmbient += OpenGl_RTVec4f (aLight.Color.r(), aLight.Color.g(), aLight.Color.b(), 0.0f);
+      myRaytraceGeometry.Ambient += BVH_Vec4f (aLight.Color.r(),
+                                               aLight.Color.g(),
+                                               aLight.Color.b(),
+                                               0.0f);
       continue;
     }
 
-    OpenGl_RTVec4f aDiffuse  (aLight.Color.r(), aLight.Color.g(), aLight.Color.b(), 1.0f);
-    OpenGl_RTVec4f aPosition (-aLight.Direction.x(), -aLight.Direction.y(), -aLight.Direction.z(), 0.0f);
+    BVH_Vec4f aDiffuse  (aLight.Color.r(),
+                         aLight.Color.g(),
+                         aLight.Color.b(),
+                         1.0f);
+
+    BVH_Vec4f aPosition (-aLight.Direction.x(),
+                         -aLight.Direction.y(),
+                         -aLight.Direction.z(),
+                         0.0f);
+
     if (aLight.Type != Visual3d_TOLS_DIRECTIONAL)
     {
-      aPosition = OpenGl_RTVec4f (aLight.Position.x(), aLight.Position.y(), aLight.Position.z(), 1.0f);
+      aPosition = BVH_Vec4f (aLight.Position.x(),
+                             aLight.Position.y(),
+                             aLight.Position.z(),
+                             1.0f);
     }
+
     if (aLight.IsHeadlight)
-    {
       aPosition = MatVecMult (theInvModelView, aPosition);
-    }
 
-    myRaytraceSceneData.LightSources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition));
+    
+    myRaytraceGeometry.Sources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition));
   }
 
-  if (myRaytraceSceneData.LightSources.size() > 0)
+  if (myRaytraceLightSrcTexture.IsNull())  // create light source buffer
   {
-    myRaytraceSceneData.LightSources.front().Ambient += anAmbient;
+    myRaytraceLightSrcTexture = new OpenGl_TextureBufferArb;
+
+    if (!myRaytraceLightSrcTexture->Create (myGlContext))
+    {
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error: Failed to create light source buffer" << std::endl;
+#endif
+      return Standard_False;
+    }
   }
-  else
+  
+  if (myRaytraceGeometry.Sources.size() != 0)
   {
-    myRaytraceSceneData.LightSources.push_back (OpenGl_RaytraceLight (OpenGl_RTVec4f (anAmbient.rgb(), -1.0f)));
+    const GLfloat* aDataPtr = myRaytraceGeometry.Sources.front().Packed();
+    if (!myRaytraceLightSrcTexture->Init (myGlContext, 4, GLsizei (myRaytraceGeometry.Sources.size() * 2), aDataPtr))
+    {
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error: Failed to upload light source buffer" << std::endl;
+#endif
+      return Standard_False;
+    }
   }
 
-  cl_int anError = CL_SUCCESS;
-
-  if (myRaytraceLightSourceBuffer != NULL)
-    clReleaseMemObject (myRaytraceLightSourceBuffer);
+  return Standard_True;
+}
 
-  const size_t myLightBufferSize = myRaytraceSceneData.LightSources.size() > 0
-                                 ? myRaytraceSceneData.LightSources.size()
-                                 : 1;
+// =======================================================================
+// function : UpdateRaytraceEnvironmentMap
+// purpose  : Updates environment map for ray-tracing
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::UpdateRaytraceEnvironmentMap()
+{
+  if (myView.IsNull())
+    return Standard_False;
 
-  myRaytraceLightSourceBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
-                                                myLightBufferSize * sizeof(OpenGl_RaytraceLight),
-                                                NULL, &anError);
+  if (myViewModificationStatus == myView->ModificationState())
+    return Standard_True;
 
-  if (myRaytraceSceneData.LightSources.size() > 0)
+  for (Standard_Integer anIdx = 0; anIdx < 2; ++anIdx)
   {
-    const void* aDataPtr = myRaytraceSceneData.LightSources.front().Packed();
-    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0,
-                                     myLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr,
-                                     0, NULL, NULL);
-  }
+    const Handle(OpenGl_ShaderProgram)& aProgram =
+      anIdx == 0 ? myRaytraceProgram : myPostFSAAProgram;
 
-#ifdef RAY_TRACE_PRINT_INFO
-  if (anError != CL_SUCCESS)
-  {
-    std::cout << "Error! Failed to set light sources!";
+    if (!aProgram.IsNull())
+    {
+      aProgram->Bind (myGlContext);
 
-    return Standard_False;
+      if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE)
+      {
+        myView->TextureEnv()->Bind (
+          myGlContext, GL_TEXTURE0 + OpenGl_RT_EnvironmentMapTexture);
+
+        aProgram->SetUniform (myGlContext, "uEnvironmentEnable", 1);
+      }
+      else
+      {
+        aProgram->SetUniform (myGlContext, "uEnvironmentEnable", 0);
+      }
+
+      aProgram->SetSampler (myGlContext,
+        "uEnvironmentMapTexture", OpenGl_RT_EnvironmentMapTexture);
+    }
   }
-#endif
+
+  OpenGl_ShaderProgram::Unbind (myGlContext);
+
+  myViewModificationStatus = myView->ModificationState();
 
   return Standard_True;
 }
 
 // =======================================================================
-// function : CheckOpenCL
-// purpose  : Checks OpenCL dynamic library availability
+// function : Source
+// purpose  : Returns shader source combined with prefix
 // =======================================================================
-Standard_Boolean CheckOpenCL()
+TCollection_AsciiString OpenGl_Workspace::ShaderSource::Source() const
 {
-#if defined ( _WIN32 )
+  static const TCollection_AsciiString aVersion = "#version 140";
 
-  __try
-  {
-    cl_uint aNbPlatforms;
-    clGetPlatformIDs (0, NULL, &aNbPlatforms);
-  }
-  __except (EXCEPTION_EXECUTE_HANDLER)
+  if (myPrefix.IsEmpty())
   {
-    return Standard_False;
+    return aVersion + "\n" + mySource;
   }
 
-#endif
-
-  return Standard_True;
+  return aVersion + "\n" + myPrefix + "\n" + mySource;
 }
 
 // =======================================================================
-// function : InitOpenCL
-// purpose  : Initializes OpenCL objects
+// function : Load
+// purpose  : Loads shader source from specified files
 // =======================================================================
-Standard_Boolean OpenGl_Workspace::InitOpenCL()
+void OpenGl_Workspace::ShaderSource::Load (
+  const TCollection_AsciiString* theFileNames, const Standard_Integer theCount)
 {
-  if (myComputeInitStatus != OpenGl_CLIS_NONE)
-  {
-    return myComputeInitStatus == OpenGl_CLIS_INIT;
-  }
+  mySource.Clear();
 
-  if (!CheckOpenCL())
+  for (Standard_Integer anIndex = 0; anIndex < theCount; ++anIndex)
   {
-    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;
-  }
+    OSD_File aFile (theFileNames[anIndex]);
 
-  // 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;
-  }
+    Standard_ASSERT_RETURN (aFile.Exists(),
+      "Error: Failed to find shader source file", /* none */);
 
-  // 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;
-    }
+    aFile.Open (OSD_ReadOnly, OSD_Protection());
 
-    if (strncmp (aName, "NVIDIA", strlen ("NVIDIA")) == 0)
-    {
-      aPrefPlatform = aPlatforms[aPlatIter];
+    TCollection_AsciiString aSource;
 
-      // Use optimizations for NVIDIA GPUs
-      myIsAmdComputePlatform = Standard_False;
-    }
-    else if (strncmp (aName, "AMD", strlen ("AMD")) == 0)
-    {
-      aPrefPlatform = (aPrefPlatform == NULL)
-                    ? aPlatforms[aPlatIter]
-                    : aPrefPlatform;
+    Standard_ASSERT_RETURN (aFile.IsOpen(),
+      "Error: Failed to open shader source file", /* none */);
+
+    aFile.Read (aSource, (Standard_Integer) aFile.Size());
 
-      // Use optimizations for ATI/AMD platform
-      myIsAmdComputePlatform = Standard_True;
+    if (!aSource.IsEmpty())
+    {
+      mySource += TCollection_AsciiString ("\n") + aSource;
     }
-  }
 
-  if (aPrefPlatform == NULL)
-  {
-    aPrefPlatform = aPlatforms[0];
+    aFile.Close();
   }
+}
 
-  // Obtain the list of devices available in the selected platform
-  cl_uint aNbDevices = 0;
-  anError = clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU,
-                            0, NULL, &aNbDevices);
+// =======================================================================
+// function : LoadShader
+// purpose  : Creates new shader object with specified source
+// =======================================================================
+Handle(OpenGl_ShaderObject) OpenGl_Workspace::LoadShader (const ShaderSource& theSource, GLenum theType)
+{
+  Handle(OpenGl_ShaderObject) aShader = new OpenGl_ShaderObject (theType);
 
-  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)
+  if (!aShader->Create (myGlContext))
   {
-    myComputeInitStatus = OpenGl_CLIS_FAIL;
+    const TCollection_ExtendedString aMessage = "Error: Failed to create shader object";
+      
     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;
-  }
+      GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
 
-  // 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;
-  }
+    aShader->Release (myGlContext.operator->());
 
-  // 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;
+    return Handle(OpenGl_ShaderObject)();
   }
 
-  anError = clBuildProgram (myRaytraceProgram, 0,
-                            NULL, NULL, NULL, NULL);
+  if (!aShader->LoadSource (myGlContext, theSource.Source()))
   {
-    // Fetch build log
-    size_t aLogLen = 0;
-    cl_int aResult = clGetProgramBuildInfo (myRaytraceProgram, aDevice,
-                                            CL_PROGRAM_BUILD_LOG, 0, NULL, &aLogLen);
+    const TCollection_ExtendedString aMessage = "Error: Failed to set shader source";
+      
+    myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+      GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
 
-    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
-      }
-    }
-  }
+    aShader->Release (myGlContext.operator->());
 
-  if (anError != CL_SUCCESS)
-  {
-    return Standard_False;
+    return Handle(OpenGl_ShaderObject)();
   }
 
-  // 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;
-  }
+  TCollection_AsciiString aBuildLog;
 
-  myRaytraceSmoothKernel = clCreateKernel (myRaytraceProgram, "MainAntialiased", &anError);
-  if (anError != CL_SUCCESS)
+  if (!aShader->Compile (myGlContext))
   {
-    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;
-  }
+    if (aShader->FetchInfoLog (myGlContext, aBuildLog))
+    {
+      const TCollection_ExtendedString aMessage =
+        TCollection_ExtendedString ("Error: Failed to compile shader object:\n") + aBuildLog;
 
-  // Create OpenCL command queue
-  // Note: For profiling set CL_QUEUE_PROFILING_ENABLE
-  cl_command_queue_properties aProps = CL_QUEUE_PROFILING_ENABLE;
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << aBuildLog << std::endl;
+#endif
 
-  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!");
+      myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+        GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+    }
+    
+    aShader->Release (myGlContext.operator->());
 
-    return Standard_False;
+    return Handle(OpenGl_ShaderObject)();
   }
 
-  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)
+#ifdef RAY_TRACE_PRINT_INFO
+  if (aShader->FetchInfoLog (myGlContext, aBuildLog))
   {
-    return Standard_False;
-  }
+    if (!aBuildLog.IsEmpty())
+    {
+      std::cout << aBuildLog << std::endl;
+    }
+    else
+    {
+      std::cout << "Info: shader build log is empty" << std::endl;
+    }
+  }  
+#endif
 
-  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;
-  }
+  return aShader;
+}
 
-  char aDeviceName[256];
-  anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_NAME, sizeof(aDeviceName), aDeviceName, NULL);
-  theInfo.Bind ("Name", aDeviceName);
+// =======================================================================
+// function : SafeFailBack
+// purpose  : Performs safe exit when shaders initialization fails
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::SafeFailBack (const TCollection_ExtendedString& theMessage)
+{
+  myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+    GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, theMessage);
 
-  char aDeviceVendor[256];
-  anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_VENDOR, sizeof(aDeviceVendor), aDeviceVendor, NULL);
-  theInfo.Bind ("Vendor", aDeviceVendor);
+  myComputeInitStatus = OpenGl_RT_FAIL;
 
-  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;
+  ReleaseRaytraceResources();
+  
+  return Standard_False;
 }
 
 // =======================================================================
-// function : ReleaseOpenCL
-// purpose  : Releases resources of OpenCL objects
+// function : InitRaytraceResources
+// purpose  : Initializes OpenGL/GLSL shader programs
 // =======================================================================
-void OpenGl_Workspace::ReleaseOpenCL()
+Standard_Boolean OpenGl_Workspace::InitRaytraceResources()
 {
-  clReleaseKernel (myRaytraceRenderKernel);
-  clReleaseKernel (myRaytraceSmoothKernel);
+  Standard_Boolean aToRebuildShaders = Standard_False;
 
-  clReleaseProgram (myRaytraceProgram);
-  clReleaseCommandQueue (myRaytraceQueue);
+  if (myComputeInitStatus == OpenGl_RT_INIT)
+  {
+    if (!myIsRaytraceDataValid)
+      return Standard_True;
 
-  clReleaseMemObject (myRaytraceOutputImage);
-  clReleaseMemObject (myRaytraceEnvironment);
-  clReleaseMemObject (myRaytraceOutputImageSmooth);
+    const Standard_Integer aRequiredStackSize =
+      myRaytraceGeometry.HighLevelTreeDepth() + myRaytraceGeometry.BottomLevelTreeDepth();
 
-  clReleaseMemObject (myRaytraceVertexBuffer);
-  clReleaseMemObject (myRaytraceNormalBuffer);
-  clReleaseMemObject (myRaytraceTriangleBuffer);
+    if (myTraversalStackSize < aRequiredStackSize)
+    {
+      myTraversalStackSize = Max (aRequiredStackSize, THE_DEFAULT_STACK_SIZE);
 
-  clReleaseMemObject (myRaytraceMaterialBuffer);
-  clReleaseMemObject (myRaytraceLightSourceBuffer);
+      aToRebuildShaders = Standard_True;
+    }
+    else
+    {
+      if (aRequiredStackSize < myTraversalStackSize)
+      {
+        if (myTraversalStackSize > THE_DEFAULT_STACK_SIZE)
+        {
+          myTraversalStackSize = Max (aRequiredStackSize, THE_DEFAULT_STACK_SIZE);
 
-  clReleaseMemObject (myRaytraceNodeMinPointBuffer);
-  clReleaseMemObject (myRaytraceNodeMaxPointBuffer);
-  clReleaseMemObject (myRaytraceNodeDataRcrdBuffer);
+          aToRebuildShaders = Standard_True;
+        }
+      }
+    }
 
-  clReleaseContext (myComputeContext);
+    if (aToRebuildShaders)
+    {
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Info: Rebuild shaders with stack size: " << myTraversalStackSize << std::endl;
+#endif
 
-  if (glIsTexture (*myRaytraceOutputTexture))
-    glDeleteTextures (2, myRaytraceOutputTexture);
-}
+      TCollection_AsciiString aStackSizeStr =
+        TCollection_AsciiString ("#define STACK_SIZE ") + TCollection_AsciiString (myTraversalStackSize);
 
-// =======================================================================
-// 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;
+      myRaytraceShaderSource.SetPrefix (aStackSizeStr);
+      myPostFSAAShaderSource.SetPrefix (aStackSizeStr);
+
+      if (!myRaytraceShader->LoadSource (myGlContext, myRaytraceShaderSource.Source())
+       || !myPostFSAAShader->LoadSource (myGlContext, myPostFSAAShaderSource.Source()))
+      {
+        return Standard_False;
+      }
+
+      if (!myRaytraceShader->Compile (myGlContext)
+       || !myPostFSAAShader->Compile (myGlContext))
+      {
+        return Standard_False;
+      }
+
+      if (!myRaytraceProgram->Link (myGlContext)
+       || !myPostFSAAProgram->Link (myGlContext))
+      {
+        return Standard_False;
+      }
+    }
   }
 
-  bool toResize = true;
-  GLint aSizeX = -1;
-  GLint aSizeY = -1;
-  if (*myRaytraceOutputTexture != 0)
+  if (myComputeInitStatus == OpenGl_RT_NONE)
   {
-    if (!myGlContext->IsGlGreaterEqual (2, 1))
+    if (!myGlContext->IsGlGreaterEqual (3, 1))
     {
+      const TCollection_ExtendedString aMessage = "Ray-tracing requires OpenGL 3.1 and higher";
+
+      myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+        GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+
       return Standard_False;
     }
 
-    glBindTexture (GL_TEXTURE_RECTANGLE, *myRaytraceOutputTexture);
+    TCollection_AsciiString aFolder = Graphic3d_ShaderProgram::ShadersFolder();
 
-    glGetTexLevelParameteriv (GL_TEXTURE_RECTANGLE, 0, GL_TEXTURE_WIDTH,  &aSizeX);
-    glGetTexLevelParameteriv (GL_TEXTURE_RECTANGLE, 0, GL_TEXTURE_HEIGHT, &aSizeY);
+    if (aFolder.IsEmpty())
+    {
+      const TCollection_ExtendedString aMessage = "Failed to locate shaders directory";
+      
+      myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+        GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+      
+      return Standard_False;
+    }
 
-    toResize = (aSizeX != theSizeX) || (aSizeY != theSizeY);
-    if (toResize)
+    if (myIsRaytraceDataValid)
     {
-      glDeleteTextures (2, myRaytraceOutputTexture);
+      myTraversalStackSize = Max (THE_DEFAULT_STACK_SIZE,
+        myRaytraceGeometry.HighLevelTreeDepth() + myRaytraceGeometry.BottomLevelTreeDepth());
     }
-  }
-  if (!toResize)
-  {
-    return Standard_True;
-  }
 
-  glGenTextures (2, myRaytraceOutputTexture);
-  for (int aTexIter = 0; aTexIter < 2; ++aTexIter)
-  {
-    glBindTexture (GL_TEXTURE_RECTANGLE, myRaytraceOutputTexture[aTexIter]);
+    {
+      Handle(OpenGl_ShaderObject) aBasicVertShader = LoadShader (
+        ShaderSource (aFolder + "/RaytraceBase.vs"), GL_VERTEX_SHADER);
 
-    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);
+      if (aBasicVertShader.IsNull())
+      {
+        return SafeFailBack ("Failed to set vertex shader source");
+      }
 
-    glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
-    glTexParameteri (GL_TEXTURE_RECTANGLE, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
+      TCollection_AsciiString aFiles[] = { aFolder + "/RaytraceBase.fs", aFolder + "/RaytraceRender.fs" };
 
-    glTexImage2D (GL_TEXTURE_RECTANGLE, 0, GL_RGBA32F,
-                  theSizeX, theSizeY, 0,
-                  GL_RGBA, GL_FLOAT, NULL);
-  }
+      myRaytraceShaderSource.Load (aFiles, 2);
+
+      TCollection_AsciiString aStackSizeStr =
+        TCollection_AsciiString ("#define STACK_SIZE ") + TCollection_AsciiString (myTraversalStackSize);
+
+      myRaytraceShaderSource.SetPrefix (aStackSizeStr);
+
+      myRaytraceShader = LoadShader (myRaytraceShaderSource, GL_FRAGMENT_SHADER);
+
+      if (myRaytraceShader.IsNull())
+      {
+        aBasicVertShader->Release (myGlContext.operator->());
+
+        return SafeFailBack ("Failed to set ray-trace fragment shader source");
+      }
+
+      myRaytraceProgram = new OpenGl_ShaderProgram;
+
+      if (!myRaytraceProgram->Create (myGlContext))
+      {
+        aBasicVertShader->Release (myGlContext.operator->());
+
+        return SafeFailBack ("Failed to create ray-trace shader program");
+      }
+
+      if (!myRaytraceProgram->AttachShader (myGlContext, aBasicVertShader)
+       || !myRaytraceProgram->AttachShader (myGlContext, myRaytraceShader))
+      {
+        aBasicVertShader->Release (myGlContext.operator->());
+
+        return SafeFailBack ("Failed to attach ray-trace shader objects");
+      }
+
+      if (!myRaytraceProgram->Link (myGlContext))
+      {
+        TCollection_AsciiString aLinkLog;
+
+        if (myRaytraceProgram->FetchInfoLog (myGlContext, aLinkLog))
+        {
+  #ifdef RAY_TRACE_PRINT_INFO
+          std::cout << aLinkLog << std::endl;
+  #endif
+        }
+
+        return SafeFailBack ("Failed to link ray-trace shader program");
+      }
+    }
+
+    {
+      Handle(OpenGl_ShaderObject) aBasicVertShader = LoadShader (
+        ShaderSource (aFolder + "/RaytraceBase.vs"), GL_VERTEX_SHADER);
 
-  cl_int anError = CL_SUCCESS;
+      if (aBasicVertShader.IsNull())
+      {
+        return SafeFailBack ("Failed to set vertex shader source");
+      }
+
+      TCollection_AsciiString aFiles[] = { aFolder + "/RaytraceBase.fs", aFolder + "/RaytraceSmooth.fs" };
+
+      myPostFSAAShaderSource.Load (aFiles, 2);
+
+      TCollection_AsciiString aStackSizeStr =
+        TCollection_AsciiString ("#define STACK_SIZE ") + TCollection_AsciiString (myTraversalStackSize);
+
+      myPostFSAAShaderSource.SetPrefix (aStackSizeStr);
+    
+      myPostFSAAShader = LoadShader (myPostFSAAShaderSource, GL_FRAGMENT_SHADER);
+
+      if (myPostFSAAShader.IsNull())
+      {
+        aBasicVertShader->Release (myGlContext.operator->());
+
+        return SafeFailBack ("Failed to set FSAA fragment shader source");
+      }
 
-  if (myRaytraceOutputImage != NULL)
+      myPostFSAAProgram = new OpenGl_ShaderProgram;
+
+      if (!myPostFSAAProgram->Create (myGlContext))
+      {
+        aBasicVertShader->Release (myGlContext.operator->());
+
+        return SafeFailBack ("Failed to create FSAA shader program");
+      }
+
+      if (!myPostFSAAProgram->AttachShader (myGlContext, aBasicVertShader)
+       || !myPostFSAAProgram->AttachShader (myGlContext, myPostFSAAShader))
+      {
+        aBasicVertShader->Release (myGlContext.operator->());
+
+        return SafeFailBack ("Failed to attach FSAA shader objects");
+      }
+
+      if (!myPostFSAAProgram->Link (myGlContext))
+      {
+        TCollection_AsciiString aLinkLog;
+
+        if (myPostFSAAProgram->FetchInfoLog (myGlContext, aLinkLog))
+        {
+  #ifdef RAY_TRACE_PRINT_INFO
+          std::cout << aLinkLog << std::endl;
+  #endif
+        }
+      
+        return SafeFailBack ("Failed to link FSAA shader program");
+      }
+    }
+  }
+
+  if (myComputeInitStatus == OpenGl_RT_NONE || aToRebuildShaders)
   {
-    clReleaseMemObject (myRaytraceOutputImage);
+    for (Standard_Integer anIndex = 0; anIndex < 2; ++anIndex)
+    {
+      Handle(OpenGl_ShaderProgram)& aShaderProgram =
+        (anIndex == 0) ? myRaytraceProgram : myPostFSAAProgram;
+
+      aShaderProgram->Bind (myGlContext);
+
+      aShaderProgram->SetSampler (myGlContext,
+        "uSceneMinPointTexture", OpenGl_RT_SceneMinPointTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uSceneMaxPointTexture", OpenGl_RT_SceneMaxPointTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uSceneNodeInfoTexture", OpenGl_RT_SceneNodeInfoTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uObjectMinPointTexture", OpenGl_RT_ObjectMinPointTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uObjectMaxPointTexture", OpenGl_RT_ObjectMaxPointTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uObjectNodeInfoTexture", OpenGl_RT_ObjectNodeInfoTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uGeometryVertexTexture", OpenGl_RT_GeometryVertexTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uGeometryNormalTexture", OpenGl_RT_GeometryNormalTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uGeometryTriangTexture", OpenGl_RT_GeometryTriangTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uRaytraceMaterialTexture", OpenGl_RT_RaytraceMaterialTexture);
+      aShaderProgram->SetSampler (myGlContext,
+        "uRaytraceLightSrcTexture", OpenGl_RT_RaytraceLightSrcTexture);
+
+      if (anIndex == 1)
+      {
+        aShaderProgram->SetSampler (myGlContext,
+          "uFSAAInputTexture", OpenGl_RT_FSAAInputTexture);
+      }
+
+      myUniformLocations[anIndex][OpenGl_RT_aPosition] =
+        aShaderProgram->GetAttributeLocation (myGlContext, "aPosition");
+
+      myUniformLocations[anIndex][OpenGl_RT_uOriginLB] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uOriginLB");
+      myUniformLocations[anIndex][OpenGl_RT_uOriginRB] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uOriginRB");
+      myUniformLocations[anIndex][OpenGl_RT_uOriginLT] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uOriginLT");
+      myUniformLocations[anIndex][OpenGl_RT_uOriginRT] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uOriginRT");
+      myUniformLocations[anIndex][OpenGl_RT_uDirectLB] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uDirectLB");
+      myUniformLocations[anIndex][OpenGl_RT_uDirectRB] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uDirectRB");
+      myUniformLocations[anIndex][OpenGl_RT_uDirectLT] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uDirectLT");
+      myUniformLocations[anIndex][OpenGl_RT_uDirectRT] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uDirectRT");
+
+      myUniformLocations[anIndex][OpenGl_RT_uLightCount] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uLightCount");
+      myUniformLocations[anIndex][OpenGl_RT_uLightAmbnt] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uGlobalAmbient");
+
+      myUniformLocations[anIndex][OpenGl_RT_uSceneRad] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uSceneRadius");
+      myUniformLocations[anIndex][OpenGl_RT_uSceneEps] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uSceneEpsilon");
+
+      myUniformLocations[anIndex][OpenGl_RT_uShadEnabled] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uShadowsEnable");
+      myUniformLocations[anIndex][OpenGl_RT_uReflEnabled] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uReflectionsEnable");
+
+      myUniformLocations[anIndex][OpenGl_RT_uOffsetX] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uOffsetX");
+      myUniformLocations[anIndex][OpenGl_RT_uOffsetY] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uOffsetY");
+      myUniformLocations[anIndex][OpenGl_RT_uSamples] =
+        aShaderProgram->GetUniformLocation (myGlContext, "uSamples");
+    }
+
+    OpenGl_ShaderProgram::Unbind (myGlContext);
   }
-  if (myRaytraceOutputImageSmooth != NULL)
+
+  if (myComputeInitStatus != OpenGl_RT_NONE)
   {
-    clReleaseMemObject (myRaytraceOutputImageSmooth);
+    return myComputeInitStatus == OpenGl_RT_INIT;
   }
 
-  myRaytraceOutputImage = clCreateFromGLTexture2D (myComputeContext, CL_MEM_READ_WRITE,
-                                                   GL_TEXTURE_RECTANGLE, 0,
-                                                   myRaytraceOutputTexture[0], &anError);
-  if (anError != CL_SUCCESS)
+  if (myRaytraceFBO1.IsNull())
   {
-#ifdef RAY_TRACE_PRINT_INFO
-    std::cout << "Error! Failed to create output image!" << std::endl;
-#endif
-    return Standard_False;
+    myRaytraceFBO1 = new OpenGl_FrameBuffer;
   }
 
-  myRaytraceOutputImageSmooth = clCreateFromGLTexture2D (myComputeContext, CL_MEM_READ_WRITE,
-                                                         GL_TEXTURE_RECTANGLE, 0,
-                                                         myRaytraceOutputTexture[1], &anError);
-  if (anError != CL_SUCCESS)
+  if (myRaytraceFBO2.IsNull())
   {
-#ifdef RAY_TRACE_PRINT_INFO
-    std::cout << "Error! Failed to create anti-aliased output image!" << std::endl;
-#endif
-    return Standard_False;
+    myRaytraceFBO2 = new OpenGl_FrameBuffer;
   }
 
+  const GLfloat aVertices[] = { -1.f, -1.f,  0.f,
+                                -1.f,  1.f,  0.f,
+                                 1.f,  1.f,  0.f,
+                                 1.f,  1.f,  0.f,
+                                 1.f, -1.f,  0.f,
+                                -1.f, -1.f,  0.f };
+
+  myRaytraceScreenQuad.Init (myGlContext, 3, 6, aVertices);
+
+  myComputeInitStatus = OpenGl_RT_INIT; // initialized in normal way
+  
   return Standard_True;
 }
 
 // =======================================================================
-// function : WriteRaytraceSceneToDevice
-// purpose  : Writes scene geometry to OpenCl device
+// function : NullifyResource
+// purpose  :
 // =======================================================================
-Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice()
+inline void NullifyResource (const Handle(OpenGl_Context)& theContext,
+                             Handle(OpenGl_Resource)&      theResource)
 {
-  if (myComputeContext == NULL)
-    return Standard_False;
+  if (!theResource.IsNull())
+  {
+    theResource->Release (theContext.operator->());
+    theResource.Nullify();
+  }
+}
 
-  cl_int anError = CL_SUCCESS;
+// =======================================================================
+// function : ReleaseRaytraceResources
+// purpose  : Releases OpenGL/GLSL shader programs
+// =======================================================================
+void OpenGl_Workspace::ReleaseRaytraceResources()
+{
+  NullifyResource (myGlContext, myRaytraceFBO1);
+  NullifyResource (myGlContext, myRaytraceFBO2);
 
-  if (myRaytraceNormalBuffer != NULL)
-    anError |= clReleaseMemObject (myRaytraceNormalBuffer);
+  NullifyResource (myGlContext, myRaytraceShader);
+  NullifyResource (myGlContext, myPostFSAAShader);
 
-  if (myRaytraceVertexBuffer != NULL)
-    anError |= clReleaseMemObject (myRaytraceVertexBuffer);
+  NullifyResource (myGlContext, myRaytraceProgram);
+  NullifyResource (myGlContext, myPostFSAAProgram);
 
-  if (myRaytraceTriangleBuffer != NULL)
-    anError |= clReleaseMemObject (myRaytraceTriangleBuffer);
+  NullifyResource (myGlContext, mySceneNodeInfoTexture);
+  NullifyResource (myGlContext, mySceneMinPointTexture);
+  NullifyResource (myGlContext, mySceneMaxPointTexture);
 
-  if (myRaytraceNodeMinPointBuffer != NULL)
-    anError |= clReleaseMemObject (myRaytraceNodeMinPointBuffer);
+  NullifyResource (myGlContext, myObjectNodeInfoTexture);
+  NullifyResource (myGlContext, myObjectMinPointTexture);
+  NullifyResource (myGlContext, myObjectMaxPointTexture);
 
-  if (myRaytraceNodeMaxPointBuffer != NULL)
-    anError |= clReleaseMemObject (myRaytraceNodeMaxPointBuffer);
+  NullifyResource (myGlContext, myGeometryVertexTexture);
+  NullifyResource (myGlContext, myGeometryNormalTexture);
+  NullifyResource (myGlContext, myGeometryTriangTexture);
 
-  if (myRaytraceNodeDataRcrdBuffer != NULL)
-    anError |= clReleaseMemObject (myRaytraceNodeDataRcrdBuffer);
+  NullifyResource (myGlContext, myRaytraceLightSrcTexture);
+  NullifyResource (myGlContext, myRaytraceMaterialTexture);
 
-  if (myRaytraceMaterialBuffer != NULL)
-    anError |= clReleaseMemObject (myRaytraceMaterialBuffer);
+  if (myRaytraceScreenQuad.IsValid())
+    myRaytraceScreenQuad.Release (myGlContext.operator->());
+}
 
-  if (anError != CL_SUCCESS)
+// =======================================================================
+// function : UploadRaytraceData
+// purpose  : Uploads ray-trace data to the GPU
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::UploadRaytraceData()
+{
+  if (!myGlContext->IsGlGreaterEqual (3, 1))
   {
 #ifdef RAY_TRACE_PRINT_INFO
-    std::cout << "Error! Failed to release OpenCL scene buffers!" << std::endl;
+    std::cout << "Error: OpenGL version is less than 3.1" << 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;
+  /////////////////////////////////////////////////////////////////////////////
+  // Create OpenGL texture buffers
 
-  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)
+  if (mySceneNodeInfoTexture.IsNull())  // create hight-level BVH buffers
   {
-#ifdef RAY_TRACE_PRINT_INFO
-    std::cout << "Error! Failed to create OpenCL geometry buffers!" << std::endl;
-#endif
-    return Standard_False;
-  }
+    mySceneNodeInfoTexture = new OpenGl_TextureBufferArb;
+    mySceneMinPointTexture = new OpenGl_TextureBufferArb;
+    mySceneMaxPointTexture = new OpenGl_TextureBufferArb;
 
-  // 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)
-  {
+    if (!mySceneNodeInfoTexture->Create (myGlContext)
+      || !mySceneMinPointTexture->Create (myGlContext)
+      || !mySceneMaxPointTexture->Create (myGlContext))
+    {
 #ifdef RAY_TRACE_PRINT_INFO
-    std::cout << "Error! Failed to create OpenCL material buffer!" << std::endl;
+      std::cout << "Error: Failed to create buffers for high-level scene BVH" << std::endl;
 #endif
-    return Standard_False;
+      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)
+  if (myObjectNodeInfoTexture.IsNull())  // create bottom-level BVH buffers
   {
-#ifdef RAY_TRACE_PRINT_INFO
-    std::cout << "Error! Failed to create OpenCL BVH buffers!" << std::endl;
-#endif
-    return Standard_False;
-  }
+    myObjectNodeInfoTexture = new OpenGl_TextureBufferArb;
+    myObjectMinPointTexture = new OpenGl_TextureBufferArb;
+    myObjectMaxPointTexture = new OpenGl_TextureBufferArb;
 
-  // 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)
+    if (!myObjectNodeInfoTexture->Create (myGlContext)
+      || !myObjectMinPointTexture->Create (myGlContext)
+      || !myObjectMaxPointTexture->Create (myGlContext))
     {
-  #ifdef RAY_TRACE_PRINT_INFO
-      std::cout << "Error! Failed to write OpenCL geometry buffers!" << std::endl;
-  #endif
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error: Failed to create buffers for bottom-level scene BVH" << 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)
+  if (myGeometryVertexTexture.IsNull())  // create geometry buffers
+  {
+    myGeometryVertexTexture = new OpenGl_TextureBufferArb;
+    myGeometryNormalTexture = new OpenGl_TextureBufferArb;
+    myGeometryTriangTexture = new OpenGl_TextureBufferArb;
+
+    if (!myGeometryVertexTexture->Create (myGlContext)
+      || !myGeometryNormalTexture->Create (myGlContext)
+      || !myGeometryTriangTexture->Create (myGlContext))
     {
-  #ifdef RAY_TRACE_PRINT_INFO
-      std::cout << "Error! Failed to write OpenCL BVH buffers!" << std::endl;
-  #endif
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error: Failed to create buffers for triangulation data" << std::endl;
+#endif
       return Standard_False;
     }
   }
 
-  // Write material buffers
-  if (myRaytraceSceneData.Materials.size() > 0)
+  if (myRaytraceMaterialTexture.IsNull())  // create material buffer
   {
-    const size_t aSize    = myRaytraceSceneData.Materials.size();
-    const void*  aDataPtr = myRaytraceSceneData.Materials.front().Packed();
+    myRaytraceMaterialTexture = new OpenGl_TextureBufferArb;
 
-    anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceMaterialBuffer, CL_FALSE,
-                                     0, aSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr,
-                                     0, NULL, NULL);
-    if (anError != CL_SUCCESS)
+    if (!myRaytraceMaterialTexture->Create (myGlContext))
     {
-  #ifdef RAY_TRACE_PRINT_INFO
-      std::cout << "Error! Failed to write OpenCL material buffer!" << std::endl;
-  #endif
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error: Failed to create buffers for material data" << 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
+  /////////////////////////////////////////////////////////////////////////////
+  // Write OpenGL texture buffers
 
-  if (anError == CL_SUCCESS)
-    myIsRaytraceDataValid = myRaytraceSceneData.Triangles.size() > 0;
+  const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = myRaytraceGeometry.BVH();
 
+  bool aResult = true;
+  if (!aBVH->NodeInfoBuffer().empty())
+  {
+    aResult &= mySceneNodeInfoTexture->Init (myGlContext, 4, GLsizei (aBVH->NodeInfoBuffer().size()),
+                                             reinterpret_cast<const GLuint*> (&aBVH->NodeInfoBuffer().front()));
+    aResult &= mySceneMinPointTexture->Init (myGlContext, 4, GLsizei (aBVH->MinPointBuffer().size()),
+                                             reinterpret_cast<const GLfloat*> (&aBVH->MinPointBuffer().front()));
+    aResult &= mySceneMaxPointTexture->Init (myGlContext, 4, GLsizei (aBVH->MaxPointBuffer().size()),
+                                             reinterpret_cast<const GLfloat*> (&aBVH->MaxPointBuffer().front()));
+  }
+  if (!aResult)
+  {
 #ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "Error: Failed to upload buffers for high-level scene BVH" << std::endl;
+#endif
+    return Standard_False;
+  }
 
-  float aMemUsed = static_cast<float> (
-    myRaytraceSceneData.Materials.size()) * sizeof (OpenGl_RaytraceMaterial);
+  Standard_Size aTotalVerticesNb = 0;
+  Standard_Size aTotalElementsNb = 0;
+  Standard_Size aTotalBVHNodesNb = 0;
 
-  aMemUsed += static_cast<float> (
-    myRaytraceSceneData.Triangles.size() * sizeof (OpenGl_RTVec4i) +
-    myRaytraceSceneData.Vertices.size()  * sizeof (OpenGl_RTVec4f) +
-    myRaytraceSceneData.Normals.size()   * sizeof (OpenGl_RTVec4f));
+  for (Standard_Integer anElemIndex = 0; anElemIndex < myRaytraceGeometry.Size(); ++anElemIndex)
+  {
+    OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
+      myRaytraceGeometry.Objects().ChangeValue (anElemIndex).operator->());
 
-  aMemUsed += static_cast<float> (
-    aTree.MinPointBuffer().size() * sizeof (OpenGl_RTVec4f) +
-    aTree.MaxPointBuffer().size() * sizeof (OpenGl_RTVec4f) +
-    aTree.DataRcrdBuffer().size() * sizeof (OpenGl_RTVec4i));
+    Standard_ASSERT_RETURN (aTriangleSet != NULL,
+      "Error: Failed to get triangulation of OpenGL element", Standard_False);
 
-  std::cout << "GPU memory used (Mb): " << aMemUsed / 1e6f << std::endl;
+    aTotalVerticesNb += aTriangleSet->Vertices.size();
+    aTotalElementsNb += aTriangleSet->Elements.size();
 
-#endif
+    Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(),
+      "Error: Failed to get bottom-level BVH of OpenGL element", Standard_False);
 
-  myRaytraceSceneData.Clear();
+    aTotalBVHNodesNb += aTriangleSet->BVH()->NodeInfoBuffer().size();
+  }
 
-  myBVHBuilder.CleanUp();
+  if (aTotalBVHNodesNb != 0)
+  {
+    aResult &= myObjectNodeInfoTexture->Init (myGlContext, 4, GLsizei (aTotalBVHNodesNb), static_cast<const GLuint*>  (NULL));
+    aResult &= myObjectMinPointTexture->Init (myGlContext, 4, GLsizei (aTotalBVHNodesNb), static_cast<const GLfloat*> (NULL));
+    aResult &= myObjectMaxPointTexture->Init (myGlContext, 4, GLsizei (aTotalBVHNodesNb), static_cast<const GLfloat*> (NULL));
+  }
 
-  return (CL_SUCCESS == anError);
-}
+  if (!aResult)
+  {
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "Error: Failed to upload buffers for bottom-level scene BVH" << std::endl;
+#endif
+    return Standard_False;
+  }
 
-#define OPENCL_GROUP_SIZE_TEST_
+  if (aTotalElementsNb != 0)
+  {
+    aResult &= myGeometryTriangTexture->Init (myGlContext, 4, GLsizei (aTotalElementsNb), static_cast<const GLuint*> (NULL));
+  }
 
-// =======================================================================
-// 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;
+  if (aTotalVerticesNb != 0)
+  {
+    aResult &= myGeometryVertexTexture->Init (myGlContext, 4, GLsizei (aTotalVerticesNb), static_cast<const GLfloat*> (NULL));
+    aResult &= myGeometryNormalTexture->Init (myGlContext, 4, GLsizei (aTotalVerticesNb), static_cast<const GLfloat*> (NULL));
+  }
 
-  ////////////////////////////////////////////////////////////
-  // 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);
+  if (!aResult)
+  {
+#ifdef RAY_TRACE_PRINT_INFO
+    std::cout << "Error: Failed to upload buffers for scene geometry" << std::endl;
+#endif
     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)
+  for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx)
+  {
+    if (!aBVH->IsOuter (aNodeIdx))
+      continue;
+
+    OpenGl_TriangleSet* aTriangleSet = myRaytraceGeometry.TriangleSet (aNodeIdx);
+
+    Standard_ASSERT_RETURN (aTriangleSet != NULL,
+      "Error: Failed to get triangulation of OpenGL element", Standard_False);
+
+    const Standard_Integer aBVHOffset = myRaytraceGeometry.AccelerationOffset (aNodeIdx);
+
+    Standard_ASSERT_RETURN (aBVHOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
+      "Error: Failed to get offset for bottom-level BVH", Standard_False);
+
+    const size_t aBVHBuffserSize = aTriangleSet->BVH()->NodeInfoBuffer().size();
+
+    if (aBVHBuffserSize != 0)
     {
-      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;
+      aResult &= myObjectNodeInfoTexture->SubData (myGlContext, aBVHOffset, GLsizei (aBVHBuffserSize),
+                                                   reinterpret_cast<const GLuint*> (&aTriangleSet->BVH()->NodeInfoBuffer().front()));
+      aResult &= myObjectMinPointTexture->SubData (myGlContext, aBVHOffset, GLsizei (aBVHBuffserSize),
+                                                   reinterpret_cast<const GLfloat*> (&aTriangleSet->BVH()->MinPointBuffer().front()));
+      aResult &= myObjectMaxPointTexture->SubData (myGlContext, aBVHOffset, GLsizei (aBVHBuffserSize),
+                                                   reinterpret_cast<const GLfloat*> (&aTriangleSet->BVH()->MaxPointBuffer().front()));
+      if (!aResult)
+      {
+#ifdef RAY_TRACE_PRINT_INFO
+        std::cout << "Error: Failed to upload buffers for bottom-level scene BVHs" << std::endl;
+#endif
+        return Standard_False;
+      }
     }
-  }
 
-  // Set work size
-  size_t aLocSizeRender[] = { myIsAmdComputePlatform ? 2 : 4, 32 };
+    const Standard_Integer aVerticesOffset = myRaytraceGeometry.VerticesOffset (aNodeIdx);
 
-#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
+    Standard_ASSERT_RETURN (aVerticesOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
+      "Error: Failed to get offset for triangulation vertices of OpenGL element", Standard_False);
 
-    size_t aWorkSizeX = theSizeX;
-    if (aWorkSizeX % aLocSizeRender[0] != 0)
-      aWorkSizeX += aLocSizeRender[0] - aWorkSizeX % aLocSizeRender[0];
+    if (!aTriangleSet->Vertices.empty())
+    {
+      aResult &= myGeometryNormalTexture->SubData (myGlContext, aVerticesOffset, GLsizei (aTriangleSet->Normals.size()),
+                                                   reinterpret_cast<const GLfloat*> (&aTriangleSet->Normals.front()));
+      aResult &= myGeometryVertexTexture->SubData (myGlContext, aVerticesOffset, GLsizei (aTriangleSet->Vertices.size()),
+                                                   reinterpret_cast<const GLfloat*> (&aTriangleSet->Vertices.front()));
+    }
 
-    size_t aWokrSizeY = theSizeY;
-    if (aWokrSizeY % aLocSizeRender[1] != 0 )
-      aWokrSizeY += aLocSizeRender[1] - aWokrSizeY % aLocSizeRender[1];
+    const Standard_Integer anElementsOffset = myRaytraceGeometry.ElementsOffset (aNodeIdx);
 
-    size_t aGlbSizeRender[] = { aWorkSizeX, aWokrSizeY };
+    Standard_ASSERT_RETURN (anElementsOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
+      "Error: Failed to get offset for triangulation elements of OpenGL element", Standard_False);
 
-    // Run kernel
-    cl_event anEvent (NULL), anEventSmooth (NULL);
-    anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceRenderKernel,
-                                      2, NULL, aGlbSizeRender, aLocSizeRender,
-                                      0, NULL, &anEvent);
-    if (anError != CL_SUCCESS)
+    if (!aTriangleSet->Elements.empty())
     {
-      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;
+      aResult &= myGeometryTriangTexture->SubData (myGlContext, anElementsOffset, GLsizei (aTriangleSet->Elements.size()),
+                                                   reinterpret_cast<const GLuint*> (&aTriangleSet->Elements.front()));
     }
-    clWaitForEvents (1, &anEvent);
 
-    if (theCView.IsAntialiasingEnabled)
+    if (!aResult)
     {
-      size_t aLocSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4,
-                                  myIsAmdComputePlatform ? 8 : 32 };
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error: Failed to upload triangulation buffers for OpenGL element" << std::endl;
+#endif
+      return Standard_False;
+    }
+  }
 
-#ifdef OPENCL_GROUP_SIZE_TEST
-      aLocSizeSmooth[0] = aLocX;
-      aLocSizeSmooth[1] = aLocY;
+  if (myRaytraceGeometry.Materials.size() != 0)
+  {
+    const GLfloat* aDataPtr = myRaytraceGeometry.Materials.front().Packed();
+    aResult &= myRaytraceMaterialTexture->Init (myGlContext, 4, GLsizei (myRaytraceGeometry.Materials.size() * 7), aDataPtr);
+    if (!aResult)
+    {
+#ifdef RAY_TRACE_PRINT_INFO
+      std::cout << "Error: Failed to upload material buffer" << std::endl;
 #endif
+      return Standard_False;
+    }
+  }
 
-      aWorkSizeX = theSizeX;
-      if (aWorkSizeX % aLocSizeSmooth[0] != 0)
-        aWorkSizeX += aLocSizeSmooth[0] - aWorkSizeX % aLocSizeSmooth[0];
+  myIsRaytraceDataValid = myRaytraceGeometry.Objects().Size() != 0;
 
-      size_t aWokrSizeY = theSizeY;
-      if (aWokrSizeY % aLocSizeSmooth[1] != 0 )
-        aWokrSizeY += aLocSizeSmooth[1] - aWokrSizeY % aLocSizeSmooth[1];
+#ifdef RAY_TRACE_PRINT_INFO
 
-      size_t aGlbSizeSmooth [] = { aWorkSizeX, aWokrSizeY };
-      anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceSmoothKernel,
-                                        2, NULL, aGlbSizeSmooth, aLocSizeSmooth,
-                                        0, NULL, &anEventSmooth);
-      clWaitForEvents (1, &anEventSmooth);
+  Standard_ShortReal aMemUsed = 0.f;
 
-      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;
-      }
-    }
+  for (Standard_Integer anElemIdx = 0; anElemIdx < myRaytraceGeometry.Size(); ++anElemIdx)
+  {
+    OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
+      myRaytraceGeometry.Objects().ChangeValue (anElemIdx).operator->());
+
+    aMemUsed += static_cast<Standard_ShortReal> (
+      aTriangleSet->Vertices.size() * sizeof (BVH_Vec4f));
+    aMemUsed += static_cast<Standard_ShortReal> (
+      aTriangleSet->Normals.size() * sizeof (BVH_Vec4f));
+    aMemUsed += static_cast<Standard_ShortReal> (
+      aTriangleSet->Elements.size() * sizeof (BVH_Vec4i));
 
-    // Get the profiling data
-#if defined (RAY_TRACE_PRINT_INFO) || defined(OPENCL_GROUP_SIZE_TEST)
+    aMemUsed += static_cast<Standard_ShortReal> (
+      aTriangleSet->BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
+    aMemUsed += static_cast<Standard_ShortReal> (
+      aTriangleSet->BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
+    aMemUsed += static_cast<Standard_ShortReal> (
+      aTriangleSet->BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
+  }
 
-    cl_ulong aTimeStart,
-             aTimeFinal;
+  aMemUsed += static_cast<Standard_ShortReal> (
+    myRaytraceGeometry.BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
+  aMemUsed += static_cast<Standard_ShortReal> (
+    myRaytraceGeometry.BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
+  aMemUsed += static_cast<Standard_ShortReal> (
+    myRaytraceGeometry.BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
 
-    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;
+  std::cout << "GPU Memory Used (MB): ~" << aMemUsed / 1048576 << 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);
+  return aResult;
+}
 
-    if (anEventSmooth != NULL)
-      clReleaseEvent (anEventSmooth);
+// =======================================================================
+// function : ResizeRaytraceBuffers
+// purpose  : Resizes OpenGL frame buffers
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::ResizeRaytraceBuffers (const Standard_Integer theSizeX,
+                                                          const Standard_Integer theSizeY)
+{
+  if (myRaytraceFBO1->GetVPSizeX() != theSizeX
+   || myRaytraceFBO1->GetVPSizeY() != theSizeY)
+  {
+    myRaytraceFBO1->Init (myGlContext, theSizeX, theSizeY);
+    myRaytraceFBO2->Init (myGlContext, theSizeX, theSizeY);
   }
 
   return Standard_True;
 }
 
 // =======================================================================
-// function : ComputeInverseMatrix
-// purpose  : Computes inversion of 4x4 floating-point matrix
+// function : UpdateCamera
+// purpose  : Generates viewing rays for corners of screen quad
 // =======================================================================
-template <typename T>
-void ComputeInverseMatrix (const T m[16], T inv[16])
+void OpenGl_Workspace::UpdateCamera (const NCollection_Mat4<GLdouble>& theOrientation,
+                                     const NCollection_Mat4<GLdouble>& theViewMapping,
+                                     OpenGl_Vec3                       theOrigins[4],
+                                     OpenGl_Vec3                       theDirects[4])
 {
-  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]);
+  NCollection_Mat4<GLdouble> aInvModelProj;
 
-  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]);
+  // compute invserse model-view-projection matrix
+  (theViewMapping * theOrientation).Inverted (aInvModelProj);
 
-  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]);
+  Standard_Integer aOriginIndex = 0;
+  Standard_Integer aDirectIndex = 0;
 
-  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]);
+  for (Standard_Integer aY = -1; aY <= 1; aY += 2)
+  {
+    for (Standard_Integer aX = -1; aX <= 1; aX += 2)
+    {
+      OpenGl_Vec4d aOrigin (GLdouble(aX),
+                            GLdouble(aY),
+                           -1.0,
+                            1.0);
 
-  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]);
+      aOrigin = aInvModelProj * aOrigin;
 
-  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]);
+      aOrigin.x() = aOrigin.x() / aOrigin.w();
+      aOrigin.y() = aOrigin.y() / aOrigin.w();
+      aOrigin.z() = aOrigin.z() / aOrigin.w();
 
-  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]);
+      OpenGl_Vec4d aDirect (GLdouble(aX),
+                            GLdouble(aY),
+                            1.0,
+                            1.0);
 
-  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]);
+      aDirect = aInvModelProj * aDirect;
 
-  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]);
+      aDirect.x() = aDirect.x() / aDirect.w();
+      aDirect.y() = aDirect.y() / aDirect.w();
+      aDirect.z() = aDirect.z() / aDirect.w();
 
-  T det = m[0] * inv[ 0] +
-          m[1] * inv[ 4] +
-          m[2] * inv[ 8] +
-          m[3] * inv[12];
+      aDirect = aDirect - aOrigin;
 
-  if (det == T (0.0)) return;
+      GLdouble aInvLen = 1.0 / sqrt (aDirect.x() * aDirect.x() +
+                                     aDirect.y() * aDirect.y() +
+                                     aDirect.z() * aDirect.z());
 
-  det = T (1.0) / det;
+      theOrigins[aOriginIndex++] = OpenGl_Vec3 (static_cast<GLfloat> (aOrigin.x()),
+                                                static_cast<GLfloat> (aOrigin.y()),
+                                                static_cast<GLfloat> (aOrigin.z()));
 
-  for (int i = 0; i < 16; ++i)
-    inv[i] *= det;
+      theDirects[aDirectIndex++] = OpenGl_Vec3 (static_cast<GLfloat> (aDirect.x() * aInvLen),
+                                                static_cast<GLfloat> (aDirect.y() * aInvLen),
+                                                static_cast<GLfloat> (aDirect.z() * aInvLen));
+    }
+  }
 }
 
 // =======================================================================
-// function : GenerateCornerRays
-// purpose  : Generates primary rays for corners of screen quad
+// function : RunRaytraceShaders
+// purpose  : Runs ray-tracing shader programs
 // =======================================================================
-void GenerateCornerRays (const GLdouble theInvModelProj[16],
-                         float          theOrigins[16],
-                         float          theDirects[16])
+Standard_Boolean OpenGl_Workspace::RunRaytraceShaders (const Graphic3d_CView& theCView,
+                                                       const Standard_Integer theSizeX,
+                                                       const Standard_Integer theSizeY,
+                                                       const OpenGl_Vec3      theOrigins[4],
+                                                       const OpenGl_Vec3      theDirects[4],
+                                                       OpenGl_FrameBuffer*    theFrameBuffer)
 {
-  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);
+  mySceneMinPointTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_SceneMinPointTexture);
+  mySceneMaxPointTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_SceneMaxPointTexture);
+  mySceneNodeInfoTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_SceneNodeInfoTexture);
+  myObjectMinPointTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_ObjectMinPointTexture);
+  myObjectMaxPointTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_ObjectMaxPointTexture);
+  myObjectNodeInfoTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_ObjectNodeInfoTexture);
+  myGeometryVertexTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_GeometryVertexTexture);
+  myGeometryNormalTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_GeometryNormalTexture);
+  myGeometryTriangTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_GeometryTriangTexture);
+  myRaytraceMaterialTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_RaytraceMaterialTexture);
+  myRaytraceLightSrcTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_RaytraceLightSrcTexture);
+
+  if (theCView.IsAntialiasingEnabled) // render source image to FBO
+  {
+    myRaytraceFBO1->BindBuffer (myGlContext);
+    
+    glDisable (GL_BLEND);
+  }
+
+  myRaytraceProgram->Bind (myGlContext);
+
+  Standard_Integer aLightSourceBufferSize =
+    static_cast<Standard_Integer> (myRaytraceGeometry.Sources.size());
+
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uOriginLB], theOrigins[0]);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uOriginRB], theOrigins[1]);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uOriginLT], theOrigins[2]);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uOriginRT], theOrigins[3]);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uDirectLB], theDirects[0]);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uDirectRB], theDirects[1]);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uDirectLT], theDirects[2]);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uDirectRT], theDirects[3]);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uSceneRad], myRaytraceSceneRadius);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uSceneEps], myRaytraceSceneEpsilon);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uLightCount], aLightSourceBufferSize);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uLightAmbnt], myRaytraceGeometry.Ambient);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uShadEnabled], theCView.IsShadowsEnabled);
+  myRaytraceProgram->SetUniform (myGlContext,
+    myUniformLocations[0][OpenGl_RT_uReflEnabled], theCView.IsReflectionsEnabled);
+
+  myGlContext->core20fwd->glEnableVertexAttribArray (
+    myUniformLocations[0][OpenGl_RT_aPosition]);
+  {
+    myGlContext->core20fwd->glVertexAttribPointer (
+      myUniformLocations[0][OpenGl_RT_aPosition], 3, GL_FLOAT, GL_FALSE, 0, NULL);
+
+    myGlContext->core15fwd->glDrawArrays (GL_TRIANGLES, 0, 6);
+  }
+  myGlContext->core20fwd->glDisableVertexAttribArray (
+    myUniformLocations[0][OpenGl_RT_aPosition]);
+  
+  if (!theCView.IsAntialiasingEnabled)
+  {
+    myRaytraceProgram->Unbind (myGlContext);
 
-      aOrigin = MatVecMult (theInvModelProj, aOrigin);
+    return Standard_True;
+  }
 
-      OpenGl_RTVec4f aDirect (float(x),
-                              float(y),
-                              1.f,
-                              1.f);
+  myGlContext->core20fwd->glActiveTexture (
+    GL_TEXTURE0 + OpenGl_RT_FSAAInputTexture); // texture unit for FBO texture
+
+  myRaytraceFBO1->BindTexture (myGlContext);
+
+  myPostFSAAProgram->Bind (myGlContext);
+
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uOriginLB], theOrigins[0]);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uOriginRB], theOrigins[1]);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uOriginLT], theOrigins[2]);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uOriginRT], theOrigins[3]);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uDirectLB], theDirects[0]);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uDirectRB], theDirects[1]);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uDirectLT], theDirects[2]);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uDirectRT], theDirects[3]);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uSceneRad], myRaytraceSceneRadius);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uSceneEps], myRaytraceSceneEpsilon);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uLightCount], aLightSourceBufferSize);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uLightAmbnt], myRaytraceGeometry.Ambient);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uShadEnabled], theCView.IsShadowsEnabled);
+  myPostFSAAProgram->SetUniform (myGlContext,
+    myUniformLocations[1][OpenGl_RT_uReflEnabled], theCView.IsReflectionsEnabled);
+
+  const Standard_ShortReal aMaxOffset = 0.559017f;
+  const Standard_ShortReal aMinOffset = 0.186339f;
+
+  myGlContext->core20fwd->glEnableVertexAttribArray (
+    myUniformLocations[1][OpenGl_RT_aPosition]);
+  
+  myGlContext->core20fwd->glVertexAttribPointer (
+    myUniformLocations[1][OpenGl_RT_aPosition], 3, GL_FLOAT, GL_FALSE, 0, NULL);
+
+  // Perform multi-pass adaptive FSAA using ping-pong technique
+  for (Standard_Integer anIt = 0; anIt < 4; ++anIt)
+  {
+    GLfloat aOffsetX = 1.f / theSizeX;
+    GLfloat aOffsetY = 1.f / theSizeY;
+
+    if (anIt < 2)
+    {
+      aOffsetX *= anIt < 1 ? aMinOffset : -aMaxOffset;
+      aOffsetY *= anIt < 1 ? aMaxOffset :  aMinOffset;
+    }
+    else
+    {
+      aOffsetX *= anIt > 2 ?  aMaxOffset : -aMinOffset;
+      aOffsetY *= anIt > 2 ? -aMinOffset : -aMaxOffset;
+    }
+    
+    myPostFSAAProgram->SetUniform (myGlContext,
+      myUniformLocations[1][OpenGl_RT_uSamples], anIt + 2);
+    myPostFSAAProgram->SetUniform (myGlContext,
+      myUniformLocations[1][OpenGl_RT_uOffsetX], aOffsetX);
+    myPostFSAAProgram->SetUniform (myGlContext,
+      myUniformLocations[1][OpenGl_RT_uOffsetY], aOffsetY);
 
-      aDirect = MatVecMult (theInvModelProj, aDirect) - aOrigin;
+    Handle(OpenGl_FrameBuffer)& aFramebuffer = anIt % 2 ? myRaytraceFBO1 : myRaytraceFBO2;
 
-      GLdouble aInvLen = 1.f / sqrt (aDirect.x() * aDirect.x() +
-                                     aDirect.y() * aDirect.y() +
-                                     aDirect.z() * aDirect.z());
+    if (anIt == 3) // disable FBO on last iteration
+    {
+      glEnable (GL_BLEND);
 
-      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;
+      if (theFrameBuffer != NULL)
+        theFrameBuffer->BindBuffer (myGlContext);
+    }
+    else
+    {
+      aFramebuffer->BindBuffer (myGlContext);
+    }
+    
+    myGlContext->core15fwd->glDrawArrays (GL_TRIANGLES, 0, 6);
 
-      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;
+    if (anIt != 3) // set input for the next pass
+    {
+      aFramebuffer->BindTexture  (myGlContext);
+      aFramebuffer->UnbindBuffer (myGlContext);
     }
   }
+
+  myGlContext->core20fwd->glDisableVertexAttribArray (
+    myUniformLocations[1][OpenGl_RT_aPosition]);
+
+  myPostFSAAProgram->Unbind (myGlContext);
+
+  return Standard_True;
 }
 
 // =======================================================================
 // function : Raytrace
-// purpose  : Redraws the window using OpenCL ray tracing
+// purpose  : Redraws the window using OpenGL/GLSL ray-tracing
 // =======================================================================
 Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView,
-                                             const int              theSizeX,
-                                             int                    theSizeY,
-                                             const Tint             theToSwap)
+                                             const Standard_Integer theSizeX,
+                                             const Standard_Integer theSizeY,
+                                             const Standard_Boolean theToSwap,
+                                             OpenGl_FrameBuffer*    theFrameBuffer)
 {
-  if (!InitOpenCL())
+  if (!UpdateRaytraceGeometry (Standard_True))
     return Standard_False;
 
-  if (!ResizeRaytraceOutputBuffer (theSizeX, theSizeY))
+  if (!InitRaytraceResources())
     return Standard_False;
 
-  if (!UpdateRaytraceEnvironmentMap())
+  if (!ResizeRaytraceBuffers (theSizeX, theSizeY))
     return Standard_False;
 
-  if (!UpdateRaytraceGeometry (Standard_True))
+  if (!UpdateRaytraceEnvironmentMap())
     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);
+  myView->GetMatrices (theOrientation, theViewMapping);
 
-  GLdouble aOrientationMatrix[16];
-  GLdouble aViewMappingMatrix[16];
-  GLdouble aOrientationInvers[16];
+  NCollection_Mat4<GLdouble> aOrientationMatrix;
+  NCollection_Mat4<GLdouble> aViewMappingMatrix;
 
-  for (int j = 0; j < 4; ++j)
-    for (int i = 0; i < 4; ++i)
+  for (Standard_Integer j = 0; j < 4; ++j)
+  {
+    for (Standard_Integer i = 0; i < 4; ++i)
     {
       aOrientationMatrix [4 * j + i] = theOrientation (i, j);
       aViewMappingMatrix [4 * j + i] = theViewMapping (i, j);
     }
+  }
+  
+  NCollection_Mat4<GLdouble> aInvOrientationMatrix;
+  aOrientationMatrix.Inverted (aInvOrientationMatrix);
 
-  ComputeInverseMatrix (aOrientationMatrix, aOrientationInvers);
-
-  if (!UpdateRaytraceLightSources (aOrientationInvers))
+  if (!UpdateRaytraceLightSources (aInvOrientationMatrix))
     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);
-  }
+  OpenGl_Vec3 aOrigins[4];
+  OpenGl_Vec3 aDirects[4];
 
-  anError |= clEnqueueReleaseGLObjects (myRaytraceQueue,
-                                        2, anImages,
-                                        0, NULL, NULL);
-  clFinish (myRaytraceQueue);
+  UpdateCamera (aOrientationMatrix,
+                aViewMappingMatrix,
+                aOrigins,
+                aDirects);
 
   // Draw background
   glPushAttrib (GL_ENABLE_BIT |
@@ -2101,41 +2000,38 @@ Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView,
 
   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);
+  if (theFrameBuffer != NULL)
+    theFrameBuffer->BindBuffer (myGlContext);
 
-  glEnable (GL_TEXTURE_RECTANGLE);
+  myView->DrawBackground (*this);
 
+  // Generate ray-traced image
   glMatrixMode (GL_PROJECTION);
   glLoadIdentity();
 
   glMatrixMode (GL_MODELVIEW);
   glLoadIdentity();
 
-  glColor3f (1.0f, 1.0f, 1.0f);
-
-  glBindTexture (GL_TEXTURE_RECTANGLE, myRaytraceOutputTexture[theCView.IsAntialiasingEnabled ? 1 : 0]);
+  glEnable (GL_BLEND);
+  glBlendFunc (GL_ONE, GL_SRC_ALPHA);
 
   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();
+    myRaytraceScreenQuad.Bind (myGlContext);
+
+    RunRaytraceShaders (theCView,
+                        theSizeX,
+                        theSizeY,
+                        aOrigins,
+                        aDirects,
+                        theFrameBuffer);
+
+    myRaytraceScreenQuad.Unbind (myGlContext);
   }
 
+  if (theFrameBuffer != NULL)
+    theFrameBuffer->UnbindBuffer (myGlContext);
+
   glPopAttrib();
 
   // Swap the buffers
@@ -2145,9 +2041,9 @@ Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView,
     myBackBufferRestored = Standard_False;
   }
   else
+  {
     glFlush();
+  }
 
   return Standard_True;
 }
-
-#endif