src/OpenGl/OpenGl_Workspace_Raytrace.cxx

   1 // Created on: 2013-08-27
   2 // Created by: Denis BOGOLEPOV
   3 // Copyright (c) 2013 OPEN CASCADE SAS
   4 //
   5 // This file is part of Open CASCADE Technology software library.
   6 //
   7 // This library is free software; you can redistribute it and / or modify it
   8 // under the terms of the GNU Lesser General Public version 2.1 as published
   9 // by the Free Software Foundation, with special exception defined in the file
  10 // OCCT_LGPL_EXCEPTION.txt. Consult the file LICENSE_LGPL_21.txt included in OCCT
  11 // distribution for complete text of the license and disclaimer of any warranty.
  12 //
  13 // Alternatively, this file may be used under the terms of Open CASCADE
  14 // commercial license or contractual agreement.
  15
  16 #ifdef HAVE_CONFIG_H
  17   #include <config.h>
  18 #endif
  19
  20 #ifdef HAVE_OPENCL
  21
  22 #include <OpenGl_Cl.hxx>
  23
  24 #if defined(_WIN32)
  25
  26   #include <windows.h>
  27   #include <wingdi.h>
  28
  29   #pragma comment (lib, "DelayImp.lib")
  30   #pragma comment (lib, "OpenCL.lib")
  31
  32 #elif defined(__APPLE__) && !defined(MACOSX_USE_GLX)
  33   #include <OpenGL/CGLCurrent.h>
  34 #else
  35   #include <GL/glx.h>
  36 #endif
  37
  38 #include <OpenGl_Context.hxx>
  39 #include <OpenGl_Texture.hxx>
  40 #include <OpenGl_View.hxx>
  41 #include <OpenGl_Workspace.hxx>
  42 #include <Standard_Assert.hxx>
  43
  44 using namespace OpenGl_Raytrace;
  45
  46 //! Use this macro to output ray-tracing debug info
  47 //#define RAY_TRACE_PRINT_INFO
  48
  49 #ifdef RAY_TRACE_PRINT_INFO
  50   #include <OSD_Timer.hxx>
  51 #endif
  52
  53 //! OpenCL source of ray-tracing kernels.
  54 extern const char THE_RAY_TRACE_OPENCL_SOURCE[];
  55
  56 // =======================================================================
  57 // function : MatVecMult
  58 // purpose  : Multiples 4x4 matrix by 4D vector
  59 // =======================================================================
  60 template<typename T>
  61 BVH_Vec4f MatVecMult (const T m[16], const BVH_Vec4f& v)
  62 {
  63   return BVH_Vec4f (
  64     static_cast<float> (m[ 0] * v.x() + m[ 4] * v.y() +
  65                         m[ 8] * v.z() + m[12] * v.w()),
  66     static_cast<float> (m[ 1] * v.x() + m[ 5] * v.y() +
  67                         m[ 9] * v.z() + m[13] * v.w()),
  68     static_cast<float> (m[ 2] * v.x() + m[ 6] * v.y() +
  69                         m[10] * v.z() + m[14] * v.w()),
  70     static_cast<float> (m[ 3] * v.x() + m[ 7] * v.y() +
  71                         m[11] * v.z() + m[15] * v.w()));
  72 }
  73
  74 // =======================================================================
  75 // function : UpdateRaytraceEnvironmentMap
  76 // purpose  : Updates environment map for ray-tracing
  77 // =======================================================================
  78 Standard_Boolean OpenGl_Workspace::UpdateRaytraceEnvironmentMap()
  79 {
  80   if (myView.IsNull())
  81     return Standard_False;
  82
  83   if (myViewModificationStatus == myView->ModificationState())
  84     return Standard_True;
  85
  86   cl_int anError = CL_SUCCESS;
  87
  88   if (myRaytraceEnvironment != NULL)
  89     clReleaseMemObject (myRaytraceEnvironment);
  90
  91   Standard_Integer aSizeX = 1;
  92   Standard_Integer aSizeY = 1;
  93
  94   if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE)
  95   {
  96     aSizeX = (myView->TextureEnv()->SizeX() <= 0) ? 1 : myView->TextureEnv()->SizeX();
  97     aSizeY = (myView->TextureEnv()->SizeY() <= 0) ? 1 : myView->TextureEnv()->SizeY();
  98   }
  99
 100   cl_image_format anImageFormat;
 101
 102   anImageFormat.image_channel_order = CL_RGBA;
 103   anImageFormat.image_channel_data_type = CL_FLOAT;
 104
 105   myRaytraceEnvironment = clCreateImage2D (myComputeContext,
 106     CL_MEM_READ_ONLY, &anImageFormat, aSizeX, aSizeY, 0, NULL, &anError);
 107
 108   cl_float* aPixelData = new cl_float[aSizeX * aSizeY * 4];
 109
 110   // Note: texture format is not compatible with OpenCL image
 111   // (it's not possible to create image directly from texture)
 112
 113   if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE)
 114   {
 115     myView->TextureEnv()->Bind (GetGlContext());
 116
 117     glGetTexImage (GL_TEXTURE_2D,
 118                    0,
 119                    GL_RGBA,
 120                    GL_FLOAT,
 121                    aPixelData);
 122
 123     myView->TextureEnv()->Unbind (GetGlContext());
 124   }
 125   else
 126   {
 127     for (Standard_Integer aPixel = 0; aPixel < aSizeX * aSizeY * 4; ++aPixel)
 128       aPixelData[aPixel] = 0.f;
 129   }
 130
 131   size_t anImageOffset[] = { 0,
 132                              0,
 133                              0 };
 134
 135   size_t anImageRegion[] = { aSizeX,
 136                              aSizeY,
 137                              1 };
 138
 139   anError |= clEnqueueWriteImage (myComputeQueue, myRaytraceEnvironment,
 140     CL_TRUE, anImageOffset, anImageRegion, 0, 0, aPixelData, 0, NULL, NULL);
 141
 142 #ifdef RAY_TRACE_PRINT_INFO
 143   if (anError != CL_SUCCESS)
 144     std::cout << "Error! Failed to write environment map image!" << std::endl;
 145 #endif
 146
 147   delete[] aPixelData;
 148
 149   myViewModificationStatus = myView->ModificationState();
 150
 151   return (anError == CL_SUCCESS);
 152 }
 153
 154 // =======================================================================
 155 // function : UpdateRaytraceGeometry
 156 // purpose  : Updates 3D scene geometry for ray tracing
 157 // =======================================================================
 158 Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theCheck)
 159 {
 160   if (myView.IsNull())
 161     return Standard_False;
 162
 163   // Note: In 'check' mode the scene geometry is analyzed for modifications
 164   // This is light-weight procedure performed for each frame
 165
 166   if (!theCheck)
 167   {
 168     myRaytraceGeometry.Clear();
 169
 170     myIsRaytraceDataValid = Standard_False;
 171   }
 172   else
 173   {
 174     if (myLayersModificationStatus != myView->LayerList().ModificationState())
 175     {
 176       return UpdateRaytraceGeometry (Standard_False);
 177     }
 178   }
 179
 180   Standard_ShortReal* aTransform (NULL);
 181
 182   // The set of processed structures (reflected to ray-tracing)
 183   // This set is used to remove out-of-date records from the
 184   // hash map of structures
 185   std::set<const OpenGl_Structure*> anElements;
 186
 187   const OpenGl_LayerList& aList = myView->LayerList();
 188
 189   for (OpenGl_SequenceOfLayers::Iterator anLayerIt (aList.Layers()); anLayerIt.More(); anLayerIt.Next())
 190   {
 191     const OpenGl_PriorityList& aPriorityList = anLayerIt.Value();
 192
 193     if (aPriorityList.NbStructures() == 0)
 194       continue;
 195
 196     const OpenGl_ArrayOfStructure& aStructArray = aPriorityList.ArrayOfStructures();
 197
 198     for (Standard_Integer anIndex = 0; anIndex < aStructArray.Length(); ++anIndex)
 199     {
 200       OpenGl_SequenceOfStructure::Iterator aStructIt;
 201
 202       for (aStructIt.Init (aStructArray (anIndex)); aStructIt.More(); aStructIt.Next())
 203       {
 204         const OpenGl_Structure* aStructure = aStructIt.Value();
 205
 206         if (theCheck)
 207         {
 208           if (CheckRaytraceStructure (aStructure))
 209           {
 210             return UpdateRaytraceGeometry (Standard_False);
 211           }
 212         }
 213         else
 214         {
 215           if (!aStructure->IsRaytracable())
 216             continue;
 217
 218           if (aStructure->Transformation()->mat != NULL)
 219           {
 220             if (aTransform == NULL)
 221               aTransform = new Standard_ShortReal[16];
 222
 223             for (Standard_Integer i = 0; i < 4; ++i)
 224               for (Standard_Integer j = 0; j < 4; ++j)
 225               {
 226                 aTransform[j * 4 + i] = aStructure->Transformation()->mat[i][j];
 227               }
 228           }
 229
 230           AddRaytraceStructure (aStructure, aTransform, anElements);
 231         }
 232       }
 233     }
 234   }
 235
 236   if (!theCheck)
 237   {
 238     // Actualize the hash map of structures -- remove out-of-date records
 239     std::map<const OpenGl_Structure*, Standard_Size>::iterator anIter = myStructureStates.begin();
 240
 241     while (anIter != myStructureStates.end())
 242     {
 243       if (anElements.find (anIter->first) == anElements.end())
 244       {
 245         myStructureStates.erase (anIter++);
 246       }
 247       else
 248       {
 249         ++anIter;
 250       }
 251     }
 252
 253     // Actualize OpenGL layer list state
 254     myLayersModificationStatus = myView->LayerList().ModificationState();
 255
 256     // Rebuild bottom-level and high-level BVHs
 257     myRaytraceGeometry.ProcessAcceleration();
 258
 259     const Standard_ShortReal aMinRadius = Max (fabs (myRaytraceGeometry.Box().CornerMin().x()), Max (
 260       fabs (myRaytraceGeometry.Box().CornerMin().y()), fabs (myRaytraceGeometry.Box().CornerMin().z())));
 261     const Standard_ShortReal aMaxRadius = Max (fabs (myRaytraceGeometry.Box().CornerMax().x()), Max (
 262       fabs (myRaytraceGeometry.Box().CornerMax().y()), fabs (myRaytraceGeometry.Box().CornerMax().z())));
 263
 264     myRaytraceSceneRadius = 2.f /* scale factor */ * Max (aMinRadius, aMaxRadius);
 265
 266     myRaytraceSceneEpsilon = Max (1e-4f,
 267       myRaytraceGeometry.Box().Size().Length() * 1e-4f);
 268
 269     return WriteRaytraceSceneToDevice();
 270   }
 271
 272   delete [] aTransform;
 273
 274   return Standard_True;
 275 }
 276
 277 // =======================================================================
 278 // function : CheckRaytraceStructure
 279 // purpose  :  Checks to see if the structure is modified
 280 // =======================================================================
 281 Standard_Boolean OpenGl_Workspace::CheckRaytraceStructure (const OpenGl_Structure* theStructure)
 282 {
 283   if (!theStructure->IsRaytracable())
 284   {
 285     // Checks to see if all ray-tracable elements were
 286     // removed from the structure
 287     if (theStructure->ModificationState() > 0)
 288     {
 289       theStructure->ResetModificationState();
 290       return Standard_True;
 291     }
 292
 293     return Standard_False;
 294   }
 295
 296   std::map<const OpenGl_Structure*, Standard_Size>::iterator aStructState = myStructureStates.find (theStructure);
 297
 298   if (aStructState != myStructureStates.end())
 299     return aStructState->second != theStructure->ModificationState();
 300
 301   return Standard_True;
 302 }
 303
 304 // =======================================================================
 305 // function : CreateMaterial
 306 // purpose  : Creates ray-tracing material properties
 307 // =======================================================================
 308 void CreateMaterial (const OPENGL_SURF_PROP& theProp, OpenGl_RaytraceMaterial& theMaterial)
 309 {
 310   const float* aSrcAmb = theProp.isphysic ? theProp.ambcol.rgb : theProp.matcol.rgb;
 311   theMaterial.Ambient = BVH_Vec4f (aSrcAmb[0] * theProp.amb,
 312                                    aSrcAmb[1] * theProp.amb,
 313                                    aSrcAmb[2] * theProp.amb,
 314                                    1.0f);
 315
 316   const float* aSrcDif = theProp.isphysic ? theProp.difcol.rgb : theProp.matcol.rgb;
 317   theMaterial.Diffuse = BVH_Vec4f (aSrcDif[0] * theProp.diff,
 318                                    aSrcDif[1] * theProp.diff,
 319                                    aSrcDif[2] * theProp.diff,
 320                                    1.0f);
 321
 322   const float aDefSpecCol[4] = {1.0f, 1.0f, 1.0f, 1.0f};
 323   const float* aSrcSpe = theProp.isphysic ? theProp.speccol.rgb : aDefSpecCol;
 324   theMaterial.Specular = BVH_Vec4f (aSrcSpe[0] * theProp.spec,
 325                                     aSrcSpe[1] * theProp.spec,
 326                                     aSrcSpe[2] * theProp.spec,
 327                                     theProp.shine);
 328
 329   const float* aSrcEms = theProp.isphysic ? theProp.emscol.rgb : theProp.matcol.rgb;
 330   theMaterial.Emission = BVH_Vec4f (aSrcEms[0] * theProp.emsv,
 331                                     aSrcEms[1] * theProp.emsv,
 332                                     aSrcEms[2] * theProp.emsv,
 333                                     1.0f);
 334
 335   // Note: Here we use sub-linear transparency function
 336   // to produce realistic-looking transparency effect
 337   theMaterial.Transparency = BVH_Vec4f (powf (theProp.trans, 0.75f),
 338                                         1.f - theProp.trans,
 339                                         1.f,
 340                                         1.f);
 341
 342   const float aMaxRefl = Max (theMaterial.Diffuse.x() + theMaterial.Specular.x(),
 343                          Max (theMaterial.Diffuse.y() + theMaterial.Specular.y(),
 344                               theMaterial.Diffuse.z() + theMaterial.Specular.z()));
 345
 346   const float aReflectionScale = 0.75f / aMaxRefl;
 347
 348   theMaterial.Reflection = BVH_Vec4f (theProp.speccol.rgb[0] * theProp.spec,
 349                                       theProp.speccol.rgb[1] * theProp.spec,
 350                                       theProp.speccol.rgb[2] * theProp.spec,
 351                                       0.f) * aReflectionScale;
 352 }
 353
 354 // =======================================================================
 355 // function : AddRaytraceStructure
 356 // purpose  : Adds OpenGL structure to ray-traced scene geometry
 357 // =======================================================================
 358 Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure* theStructure,
 359   const Standard_ShortReal* theTransform, std::set<const OpenGl_Structure*>& theElements)
 360 {
 361   theElements.insert (theStructure);
 362
 363   if (!theStructure->IsVisible())
 364   {
 365     myStructureStates[theStructure] = theStructure->ModificationState();
 366     return Standard_True;
 367   }
 368
 369   // Get structure material
 370   Standard_Integer aStructMatID = -1;
 371
 372   if (theStructure->AspectFace() != NULL)
 373   {
 374     aStructMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
 375
 376     OpenGl_RaytraceMaterial aStructMaterial;
 377     CreateMaterial (theStructure->AspectFace()->IntFront(), aStructMaterial);
 378
 379     myRaytraceGeometry.Materials.push_back (aStructMaterial);
 380   }
 381
 382   for (OpenGl_ListOfGroup::Iterator anItg (theStructure->Groups()); anItg.More(); anItg.Next())
 383   {
 384     // Get group material
 385     Standard_Integer aGroupMatID = -1;
 386
 387     if (anItg.Value()->AspectFace() != NULL)
 388     {
 389       aGroupMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
 390
 391       OpenGl_RaytraceMaterial aGroupMaterial;
 392       CreateMaterial (anItg.Value()->AspectFace()->IntFront(), aGroupMaterial);
 393
 394       myRaytraceGeometry.Materials.push_back (aGroupMaterial);
 395     }
 396
 397     Standard_Integer aMatID = aGroupMatID < 0 ? aStructMatID : aGroupMatID;
 398
 399     if (aMatID < 0)
 400     {
 401       aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
 402
 403       myRaytraceGeometry.Materials.push_back (OpenGl_RaytraceMaterial());
 404     }
 405
 406     // Add OpenGL elements from group (extract primitives arrays and aspects)
 407     for (const OpenGl_ElementNode* aNode = anItg.Value()->FirstNode(); aNode != NULL; aNode = aNode->next)
 408     {
 409       OpenGl_AspectFace* anAspect = dynamic_cast<OpenGl_AspectFace*> (aNode->elem);
 410       if (anAspect != NULL)
 411       {
 412         aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
 413
 414         OpenGl_RaytraceMaterial aMaterial;
 415         CreateMaterial (anAspect->IntFront(), aMaterial);
 416
 417         myRaytraceGeometry.Materials.push_back (aMaterial);
 418       }
 419       else
 420       {
 421         OpenGl_PrimitiveArray* aPrimArray = dynamic_cast<OpenGl_PrimitiveArray*> (aNode->elem);
 422         if (aPrimArray != NULL)
 423         {
 424           NCollection_Handle<BVH_Object<Standard_ShortReal, 4> > aSet =
 425             AddRaytracePrimitiveArray (aPrimArray->PArray(), aMatID, theTransform);
 426
 427           if (!aSet.IsNull())
 428             myRaytraceGeometry.Objects().Append (aSet);
 429         }
 430       }
 431     }
 432   }
 433
 434   Standard_ShortReal* aTransform (NULL);
 435
 436   // Process all connected OpenGL structures
 437   for (OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures()); anIts.More(); anIts.Next())
 438   {
 439     if (anIts.Value()->Transformation()->mat != NULL)
 440     {
 441       Standard_ShortReal* aTransform = new Standard_ShortReal[16];
 442
 443       for (Standard_Integer i = 0; i < 4; ++i)
 444         for (Standard_Integer j = 0; j < 4; ++j)
 445         {
 446           aTransform[j * 4 + i] =
 447             anIts.Value()->Transformation()->mat[i][j];
 448         }
 449     }
 450
 451     if (anIts.Value()->IsRaytracable())
 452       AddRaytraceStructure (anIts.Value(), aTransform != NULL ? aTransform : theTransform, theElements);
 453   }
 454
 455   delete[] aTransform;
 456
 457   myStructureStates[theStructure] = theStructure->ModificationState();
 458
 459   return Standard_True;
 460 }
 461
 462 // =======================================================================
 463 // function : AddRaytracePrimitiveArray
 464 // purpose  : Adds OpenGL primitive array to ray-traced scene geometry
 465 // =======================================================================
 466 OpenGl_TriangleSet* OpenGl_Workspace::AddRaytracePrimitiveArray (
 467   const CALL_DEF_PARRAY* theArray, Standard_Integer theMatID, const Standard_ShortReal* theTransform)
 468 {
 469   if (theArray->type != TelPolygonsArrayType &&
 470       theArray->type != TelTrianglesArrayType &&
 471       theArray->type != TelQuadranglesArrayType &&
 472       theArray->type != TelTriangleFansArrayType &&
 473       theArray->type != TelTriangleStripsArrayType &&
 474       theArray->type != TelQuadrangleStripsArrayType)
 475   {
 476     return NULL;
 477   }
 478
 479   if (theArray->vertices == NULL)
 480     return NULL;
 481
 482 #ifdef RAY_TRACE_PRINT_INFO
 483   switch (theArray->type)
 484   {
 485     case TelPolygonsArrayType:
 486       std::cout << "\tAdding TelPolygonsArrayType" << std::endl; break;
 487     case TelTrianglesArrayType:
 488       std::cout << "\tAdding TelTrianglesArrayType" << std::endl; break;
 489     case TelQuadranglesArrayType:
 490       std::cout << "\tAdding TelQuadranglesArrayType" << std::endl; break;
 491     case TelTriangleFansArrayType:
 492       std::cout << "\tAdding TelTriangleFansArrayType" << std::endl; break;
 493     case TelTriangleStripsArrayType:
 494       std::cout << "\tAdding TelTriangleStripsArrayType" << std::endl; break;
 495     case TelQuadrangleStripsArrayType:
 496       std::cout << "\tAdding TelQuadrangleStripsArrayType" << std::endl; break;
 497   }
 498 #endif
 499
 500   OpenGl_TriangleSet* aSet = new OpenGl_TriangleSet;
 501
 502   {
 503     aSet->Vertices.reserve (theArray->num_vertexs);
 504
 505     for (Standard_Integer aVert = 0; aVert < theArray->num_vertexs; ++aVert)
 506     {
 507       BVH_Vec4f aVertex (theArray->vertices[aVert].xyz[0],
 508                          theArray->vertices[aVert].xyz[1],
 509                          theArray->vertices[aVert].xyz[2],
 510                          1.f);
 511       if (theTransform)
 512         aVertex = MatVecMult (theTransform, aVertex);
 513
 514       aSet->Vertices.push_back (aVertex);
 515     }
 516
 517     aSet->Normals.reserve (theArray->num_vertexs);
 518
 519     for (Standard_Integer aNorm = 0; aNorm < theArray->num_vertexs; ++aNorm)
 520     {
 521       BVH_Vec4f aNormal;
 522
 523       // Note: In case of absence of normals, the
 524       // renderer uses generated geometric normals
 525
 526       if (theArray->vnormals != NULL)
 527       {
 528         aNormal = BVH_Vec4f (theArray->vnormals[aNorm].xyz[0],
 529                              theArray->vnormals[aNorm].xyz[1],
 530                              theArray->vnormals[aNorm].xyz[2],
 531                              0.f);
 532
 533         if (theTransform)
 534           aNormal = MatVecMult (theTransform, aNormal);
 535       }
 536
 537       aSet->Normals.push_back (aNormal);
 538     }
 539
 540     if (theArray->num_bounds > 0)
 541     {
 542   #ifdef RAY_TRACE_PRINT_INFO
 543       std::cout << "\tNumber of bounds = " << theArray->num_bounds << std::endl;
 544   #endif
 545
 546       Standard_Integer aBoundStart = 0;
 547
 548       for (Standard_Integer aBound = 0; aBound < theArray->num_bounds; ++aBound)
 549       {
 550         const Standard_Integer aVertNum = theArray->bounds[aBound];
 551
 552   #ifdef RAY_TRACE_PRINT_INFO
 553         std::cout << "\tAdding indices from bound " << aBound << ": " <<
 554                                       aBoundStart << " .. " << aVertNum << std::endl;
 555   #endif
 556
 557         if (!AddRaytraceVertexIndices (aSet, theArray, aBoundStart, aVertNum, theMatID))
 558         {
 559           delete aSet;
 560           return NULL;
 561         }
 562
 563         aBoundStart += aVertNum;
 564       }
 565     }
 566     else
 567     {
 568       const Standard_Integer aVertNum = theArray->num_edges > 0 ? theArray->num_edges : theArray->num_vertexs;
 569
 570   #ifdef RAY_TRACE_PRINT_INFO
 571         std::cout << "\tAdding indices from array: " << aVertNum << std::endl;
 572   #endif
 573
 574       if (!AddRaytraceVertexIndices (aSet, theArray, 0, aVertNum, theMatID))
 575       {
 576         delete aSet;
 577         return NULL;
 578       }
 579     }
 580   }
 581
 582   if (aSet->Size() != 0)
 583     aSet->MarkDirty();
 584
 585   return aSet;
 586 }
 587
 588 // =======================================================================
 589 // function : AddRaytraceVertexIndices
 590 // purpose  : Adds vertex indices to ray-traced scene geometry
 591 // =======================================================================
 592 Standard_Boolean OpenGl_Workspace::AddRaytraceVertexIndices (OpenGl_TriangleSet* theSet,
 593   const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 594 {
 595   switch (theArray->type)
 596   {
 597     case TelTrianglesArrayType:
 598       return AddRaytraceTriangleArray (theSet, theArray, theOffset, theCount, theMatID);
 599
 600     case TelQuadranglesArrayType:
 601       return AddRaytraceQuadrangleArray (theSet, theArray, theOffset, theCount, theMatID);
 602
 603     case TelTriangleFansArrayType:
 604       return AddRaytraceTriangleFanArray (theSet, theArray, theOffset, theCount, theMatID);
 605
 606     case TelTriangleStripsArrayType:
 607       return AddRaytraceTriangleStripArray (theSet, theArray, theOffset, theCount, theMatID);
 608
 609     case TelQuadrangleStripsArrayType:
 610       return AddRaytraceQuadrangleStripArray (theSet, theArray, theOffset, theCount, theMatID);
 611
 612     default:
 613       return AddRaytracePolygonArray (theSet, theArray, theOffset, theCount, theMatID);
 614   }
 615 }
 616
 617 // =======================================================================
 618 // function : AddRaytraceTriangleArray
 619 // purpose  : Adds OpenGL triangle array to ray-traced scene geometry
 620 // =======================================================================
 621 Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleArray (OpenGl_TriangleSet* theSet,
 622   const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 623 {
 624   if (theCount < 3)
 625     return Standard_True;
 626
 627   theSet->Elements.reserve (theSet->Elements.size() + theCount / 3);
 628
 629   if (theArray->num_edges > 0)
 630   {
 631     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; aVert += 3)
 632     {
 633       theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
 634                                              theArray->edges[aVert + 1],
 635                                              theArray->edges[aVert + 2],
 636                                              theMatID));
 637     }
 638   }
 639   else
 640   {
 641     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; aVert += 3)
 642     {
 643       theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
 644                                              aVert + 1,
 645                                              aVert + 2,
 646                                              theMatID));
 647     }
 648   }
 649
 650   return Standard_True;
 651 }
 652
 653 // =======================================================================
 654 // function : AddRaytraceTriangleFanArray
 655 // purpose  : Adds OpenGL triangle fan array to ray-traced scene geometry
 656 // =======================================================================
 657 Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleFanArray (OpenGl_TriangleSet* theSet,
 658   const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 659 {
 660   if (theCount < 3)
 661     return Standard_True;
 662
 663   theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
 664
 665   if (theArray->num_edges > 0)
 666   {
 667     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
 668     {
 669       theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset],
 670                                              theArray->edges[aVert + 1],
 671                                              theArray->edges[aVert + 2],
 672                                              theMatID));
 673     }
 674   }
 675   else
 676   {
 677     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
 678     {
 679       theSet->Elements.push_back (BVH_Vec4i (theOffset,
 680                                              aVert + 1,
 681                                              aVert + 2,
 682                                              theMatID));
 683     }
 684   }
 685
 686   return Standard_True;
 687 }
 688
 689 // =======================================================================
 690 // function : AddRaytraceTriangleStripArray
 691 // purpose  : Adds OpenGL triangle strip array to ray-traced scene geometry
 692 // =======================================================================
 693 Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleStripArray (OpenGl_TriangleSet* theSet,
 694   const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 695 {
 696   if (theCount < 3)
 697     return Standard_True;
 698
 699   theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
 700
 701   if (theArray->num_edges > 0)
 702   {
 703     for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2)
 704     {
 705       theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + aCW ? 1 : 0],
 706                                              theArray->edges[aVert + aCW ? 0 : 1],
 707                                              theArray->edges[aVert + 2],
 708                                              theMatID));
 709     }
 710   }
 711   else
 712   {
 713     for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2)
 714     {
 715       theSet->Elements.push_back (BVH_Vec4i (aVert + aCW ? 1 : 0,
 716                                              aVert + aCW ? 0 : 1,
 717                                              aVert + 2,
 718                                              theMatID));
 719     }
 720   }
 721
 722   return Standard_True;
 723 }
 724
 725 // =======================================================================
 726 // function : AddRaytraceQuadrangleArray
 727 // purpose  : Adds OpenGL quad array to ray-traced scene geometry
 728 // =======================================================================
 729 Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleArray (OpenGl_TriangleSet* theSet,
 730   const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 731 {
 732   if (theCount < 4)
 733     return Standard_True;
 734
 735   theSet->Elements.reserve (theSet->Elements.size() + theCount / 2);
 736
 737   if (theArray->num_edges > 0)
 738   {
 739     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 4)
 740     {
 741       theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
 742                                              theArray->edges[aVert + 1],
 743                                              theArray->edges[aVert + 2],
 744                                              theMatID));
 745
 746       theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
 747                                              theArray->edges[aVert + 2],
 748                                              theArray->edges[aVert + 3],
 749                                              theMatID));
 750     }
 751   }
 752   else
 753   {
 754     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 4)
 755     {
 756       theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
 757                                              aVert + 1,
 758                                              aVert + 2,
 759                                              theMatID));
 760
 761       theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
 762                                              aVert + 2,
 763                                              aVert + 3,
 764                                              theMatID));
 765     }
 766   }
 767
 768   return Standard_True;
 769 }
 770
 771 // =======================================================================
 772 // function : AddRaytraceQuadrangleStripArray
 773 // purpose  : Adds OpenGL quad strip array to ray-traced scene geometry
 774 // =======================================================================
 775 Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleStripArray (OpenGl_TriangleSet* theSet,
 776   const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 777 {
 778   if (theCount < 4)
 779     return Standard_True;
 780
 781   theSet->Elements.reserve (theSet->Elements.size() + 2 * theCount - 6);
 782
 783   if (theArray->num_edges > 0)
 784   {
 785     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 2)
 786     {
 787       theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
 788                                              theArray->edges[aVert + 1],
 789                                              theArray->edges[aVert + 2],
 790                                              theMatID));
 791
 792       theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 1],
 793                                              theArray->edges[aVert + 3],
 794                                              theArray->edges[aVert + 2],
 795                                              theMatID));
 796     }
 797   }
 798   else
 799   {
 800     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 2)
 801     {
 802       theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
 803                                              aVert + 1,
 804                                              aVert + 2,
 805                                              theMatID));
 806
 807       theSet->Elements.push_back (BVH_Vec4i (aVert + 1,
 808                                              aVert + 3,
 809                                              aVert + 2,
 810                                              theMatID));
 811     }
 812   }
 813
 814   return Standard_True;
 815 }
 816
 817 // =======================================================================
 818 // function : AddRaytracePolygonArray
 819 // purpose  : Adds OpenGL polygon array to ray-traced scene geometry
 820 // =======================================================================
 821 Standard_Boolean OpenGl_Workspace::AddRaytracePolygonArray (OpenGl_TriangleSet* theSet,
 822   const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
 823 {
 824   if (theCount < 3)
 825     return Standard_True;
 826
 827   theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
 828
 829   if (theArray->num_edges > 0)
 830   {
 831     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
 832     {
 833       theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset],
 834                                              theArray->edges[aVert + 1],
 835                                              theArray->edges[aVert + 2],
 836                                              theMatID));
 837     }
 838   }
 839   else
 840   {
 841     for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
 842     {
 843       theSet->Elements.push_back (BVH_Vec4i (theOffset,
 844                                              aVert + 1,
 845                                              aVert + 2,
 846                                              theMatID));
 847     }
 848   }
 849
 850   return Standard_True;
 851 }
 852
 853 // =======================================================================
 854 // function : UpdateRaytraceLightSources
 855 // purpose  : Updates 3D scene light sources for ray-tracing
 856 // =======================================================================
 857 Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble theInvModelView[16])
 858 {
 859   myRaytraceGeometry.Sources.clear();
 860
 861   myRaytraceGeometry.GlobalAmbient = BVH_Vec4f (0.0f, 0.0f, 0.0f, 0.0f);
 862
 863   for (OpenGl_ListOfLight::Iterator anItl (myView->LightList()); anItl.More(); anItl.Next())
 864   {
 865     const OpenGl_Light& aLight = anItl.Value();
 866
 867     if (aLight.Type == Visual3d_TOLS_AMBIENT)
 868     {
 869       myRaytraceGeometry.GlobalAmbient += BVH_Vec4f (aLight.Color.r(),
 870                                                      aLight.Color.g(),
 871                                                      aLight.Color.b(),
 872                                                      0.0f);
 873       continue;
 874     }
 875
 876     BVH_Vec4f aDiffuse  (aLight.Color.r(),
 877                          aLight.Color.g(),
 878                          aLight.Color.b(),
 879                          1.0f);
 880
 881     BVH_Vec4f aPosition (-aLight.Direction.x(),
 882                          -aLight.Direction.y(),
 883                          -aLight.Direction.z(),
 884                          0.0f);
 885
 886     if (aLight.Type != Visual3d_TOLS_DIRECTIONAL)
 887     {
 888       aPosition = BVH_Vec4f (aLight.Position.x(),
 889                              aLight.Position.y(),
 890                              aLight.Position.z(),
 891                              1.0f);
 892     }
 893
 894     if (aLight.IsHeadlight)
 895       aPosition = MatVecMult (theInvModelView, aPosition);
 896
 897     myRaytraceGeometry.Sources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition));
 898   }
 899
 900   cl_int anError = CL_SUCCESS;
 901
 902   if (myRaytraceLightSourceBuffer != NULL)
 903     clReleaseMemObject (myRaytraceLightSourceBuffer);
 904
 905   Standard_Integer aLightBufferSize = myRaytraceGeometry.Sources.size() != 0 ?
 906     static_cast<Standard_Integer> (myRaytraceGeometry.Sources.size()) : 1;
 907
 908   myRaytraceLightSourceBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
 909     aLightBufferSize * sizeof(OpenGl_RaytraceLight), NULL, &anError);
 910
 911   if (myRaytraceGeometry.Sources.size() != 0)
 912   {
 913     const void* aDataPtr = myRaytraceGeometry.Sources.front().Packed();
 914
 915     anError |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0,
 916       aLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr, 0, NULL, NULL);
 917   }
 918
 919 #ifdef RAY_TRACE_PRINT_INFO
 920   if (anError != CL_SUCCESS)
 921   {
 922     std::cout << "Error! Failed to set light sources";
 923
 924     return Standard_False;
 925   }
 926 #endif
 927
 928   return Standard_True;
 929 }
 930
 931 // =======================================================================
 932 // function : CheckOpenCL
 933 // purpose  : Checks OpenCL dynamic library availability
 934 // =======================================================================
 935 Standard_Boolean CheckOpenCL()
 936 {
 937 #if defined ( _WIN32 )
 938
 939   __try
 940   {
 941     cl_uint aNbPlatforms;
 942     clGetPlatformIDs (0, NULL, &aNbPlatforms);
 943   }
 944   __except (EXCEPTION_EXECUTE_HANDLER)
 945   {
 946     return Standard_False;
 947   }
 948
 949 #endif
 950
 951   return Standard_True;
 952 }
 953
 954 // =======================================================================
 955 // function : InitOpenCL
 956 // purpose  : Initializes OpenCL objects
 957 // =======================================================================
 958 Standard_Boolean OpenGl_Workspace::InitOpenCL()
 959 {
 960   if (myComputeInitStatus != OpenGl_CLIS_NONE)
 961   {
 962     return myComputeInitStatus == OpenGl_CLIS_INIT;
 963   }
 964
 965   if (!CheckOpenCL())
 966   {
 967     myComputeInitStatus = OpenGl_CLIS_FAIL; // fail to load OpenCL library
 968     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
 969                               GL_DEBUG_TYPE_ERROR_ARB,
 970                               0,
 971                               GL_DEBUG_SEVERITY_HIGH_ARB,
 972                               "Failed to load OpenCL dynamic library!");
 973     return Standard_False;
 974   }
 975
 976   // Obtain the list of platforms available
 977   cl_uint aNbPlatforms = 0;
 978   cl_int  anError = clGetPlatformIDs (0, NULL, &aNbPlatforms);
 979   cl_platform_id* aPlatforms = (cl_platform_id* )alloca (aNbPlatforms * sizeof(cl_platform_id));
 980   anError |= clGetPlatformIDs (aNbPlatforms, aPlatforms, NULL);
 981   if (anError != CL_SUCCESS
 982    || aNbPlatforms == 0)
 983   {
 984     myComputeInitStatus = OpenGl_CLIS_FAIL;
 985     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
 986                               GL_DEBUG_TYPE_ERROR_ARB,
 987                               0,
 988                               GL_DEBUG_SEVERITY_HIGH_ARB,
 989                               "No any OpenCL platform installed!");
 990     return Standard_False;
 991   }
 992
 993   // Note: We try to find NVIDIA or AMD platforms with GPU devices!
 994   cl_platform_id aPrefPlatform = NULL;
 995   for (cl_uint aPlatIter = 0; aPlatIter < aNbPlatforms; ++aPlatIter)
 996   {
 997     char aName[256];
 998     anError = clGetPlatformInfo (aPlatforms[aPlatIter], CL_PLATFORM_NAME,
 999                                  sizeof(aName), aName, NULL);
1000     if (anError != CL_SUCCESS)
1001     {
1002       continue;
1003     }
1004
1005     if (strncmp (aName, "NVIDIA", strlen ("NVIDIA")) == 0)
1006     {
1007       aPrefPlatform = aPlatforms[aPlatIter];
1008
1009       // Use optimizations for NVIDIA GPUs
1010       myIsAmdComputePlatform = Standard_False;
1011     }
1012     else if (strncmp (aName, "AMD", strlen ("AMD")) == 0)
1013     {
1014       aPrefPlatform = (aPrefPlatform == NULL)
1015                     ? aPlatforms[aPlatIter]
1016                     : aPrefPlatform;
1017
1018       // Use optimizations for ATI/AMD platform
1019       myIsAmdComputePlatform = Standard_True;
1020     }
1021   }
1022
1023   if (aPrefPlatform == NULL)
1024   {
1025     aPrefPlatform = aPlatforms[0];
1026   }
1027
1028   // Obtain the list of devices available in the selected platform
1029   cl_uint aNbDevices = 0;
1030   anError = clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU,
1031                             0, NULL, &aNbDevices);
1032
1033   cl_device_id* aDevices = (cl_device_id* )alloca (aNbDevices * sizeof(cl_device_id));
1034   anError |= clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU,
1035                              aNbDevices, aDevices, NULL);
1036   if (anError != CL_SUCCESS)
1037   {
1038     myComputeInitStatus = OpenGl_CLIS_FAIL;
1039     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1040                               GL_DEBUG_TYPE_ERROR_ARB,
1041                               0,
1042                               GL_DEBUG_SEVERITY_HIGH_ARB,
1043                               "Failed to get OpenCL GPU device!");
1044     return Standard_False;
1045   }
1046
1047   // Note: Simply get first available GPU
1048   cl_device_id aDevice = aDevices[0];
1049
1050   // detect old contexts
1051   char aVerClStr[256];
1052   clGetDeviceInfo (aDevice, CL_DEVICE_VERSION,
1053                    sizeof(aVerClStr), aVerClStr, NULL);
1054   aVerClStr[strlen ("OpenCL 1.0")] = '\0';
1055   const bool isVer10 = strncmp (aVerClStr, "OpenCL 1.0", strlen ("OpenCL 1.0")) == 0;
1056
1057   // Create OpenCL context
1058   cl_context_properties aCtxProp[] =
1059   {
1060   #if defined(__APPLE__) && !defined(MACOSX_USE_GLX)
1061     CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE,
1062     (cl_context_properties )CGLGetShareGroup (CGLGetCurrentContext()),
1063   #elif defined(_WIN32)
1064     CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform,
1065     CL_GL_CONTEXT_KHR,   (cl_context_properties )wglGetCurrentContext(),
1066     CL_WGL_HDC_KHR,      (cl_context_properties )wglGetCurrentDC(),
1067   #else
1068     CL_GL_CONTEXT_KHR,   (cl_context_properties )glXGetCurrentContext(),
1069     CL_GLX_DISPLAY_KHR,  (cl_context_properties )glXGetCurrentDisplay(),
1070     CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform,
1071   #endif
1072     0
1073   };
1074
1075   myComputeContext = clCreateContext (aCtxProp,
1076                                     #if defined(__APPLE__) && !defined(MACOSX_USE_GLX)
1077                                       0, NULL, // device will be taken from GL context
1078                                     #else
1079                                       1, &aDevice,
1080                                     #endif
1081                                       NULL, NULL, &anError);
1082   if (anError != CL_SUCCESS)
1083   {
1084     myComputeInitStatus = OpenGl_CLIS_FAIL;
1085     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1086                               GL_DEBUG_TYPE_ERROR_ARB,
1087                               0,
1088                               GL_DEBUG_SEVERITY_HIGH_ARB,
1089                               "Failed to initialize OpenCL context!");
1090     return Standard_False;
1091   }
1092
1093   // Create OpenCL program
1094   const char* aSources[] =
1095   {
1096     isVer10 ? "#define M_PI_F ( float )( 3.14159265359f )\n" : "",
1097     THE_RAY_TRACE_OPENCL_SOURCE
1098   };
1099   myRaytraceProgram = clCreateProgramWithSource (myComputeContext, 2,
1100                                                  aSources, NULL, &anError);
1101   if (anError != CL_SUCCESS)
1102   {
1103     myComputeInitStatus = OpenGl_CLIS_FAIL;
1104     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1105                               GL_DEBUG_TYPE_ERROR_ARB,
1106                               0,
1107                               GL_DEBUG_SEVERITY_HIGH_ARB,
1108                               "Failed to create OpenCL ray-tracing program!");
1109     return Standard_False;
1110   }
1111
1112   anError = clBuildProgram (myRaytraceProgram, 0,
1113                             NULL, NULL, NULL, NULL);
1114   {
1115     // Fetch build log
1116     size_t aLogLen = 0;
1117     cl_int aResult = clGetProgramBuildInfo (myRaytraceProgram, aDevice,
1118                                             CL_PROGRAM_BUILD_LOG, 0, NULL, &aLogLen);
1119
1120     char* aBuildLog = (char* )alloca (aLogLen);
1121     aResult |= clGetProgramBuildInfo (myRaytraceProgram, aDevice,
1122                                       CL_PROGRAM_BUILD_LOG, aLogLen, aBuildLog, NULL);
1123     if (aResult == CL_SUCCESS)
1124     {
1125       if (anError != CL_SUCCESS)
1126       {
1127         myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1128                                   GL_DEBUG_TYPE_ERROR_ARB,
1129                                   0,
1130                                   GL_DEBUG_SEVERITY_HIGH_ARB,
1131                                   aBuildLog);
1132       }
1133       else
1134       {
1135       #ifdef RAY_TRACE_PRINT_INFO
1136         std::cout << aBuildLog << std::endl;
1137       #endif
1138       }
1139     }
1140   }
1141
1142   if (anError != CL_SUCCESS)
1143   {
1144     return Standard_False;
1145   }
1146
1147   // Create OpenCL ray tracing kernels
1148   myRaytraceRenderKernel = clCreateKernel (myRaytraceProgram, "RaytraceRender", &anError);
1149   if (anError != CL_SUCCESS)
1150   {
1151     myComputeInitStatus = OpenGl_CLIS_FAIL;
1152     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1153                               GL_DEBUG_TYPE_ERROR_ARB,
1154                               0,
1155                               GL_DEBUG_SEVERITY_HIGH_ARB,
1156                               "Failed to create OpenCL ray-tracing kernel!");
1157     return Standard_False;
1158   }
1159
1160   myRaytraceSmoothKernel = clCreateKernel (myRaytraceProgram, "RaytraceSmooth", &anError);
1161   if (anError != CL_SUCCESS)
1162   {
1163     myComputeInitStatus = OpenGl_CLIS_FAIL;
1164     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1165                               GL_DEBUG_TYPE_ERROR_ARB,
1166                               0,
1167                               GL_DEBUG_SEVERITY_HIGH_ARB,
1168                               "Failed to create OpenCL ray-tracing kernel!");
1169     return Standard_False;
1170   }
1171
1172   // Create OpenCL command queue
1173   // Note: For profiling set CL_QUEUE_PROFILING_ENABLE
1174   cl_command_queue_properties aProps = CL_QUEUE_PROFILING_ENABLE;
1175
1176   myComputeQueue = clCreateCommandQueue (myComputeContext, aDevice, aProps, &anError);
1177   if (anError != CL_SUCCESS)
1178   {
1179     myComputeInitStatus = OpenGl_CLIS_FAIL;
1180     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1181                               GL_DEBUG_TYPE_ERROR_ARB,
1182                               0,
1183                               GL_DEBUG_SEVERITY_HIGH_ARB,
1184                               "Failed to create OpenCL command queue!");
1185
1186     return Standard_False;
1187   }
1188
1189   myComputeInitStatus = OpenGl_CLIS_INIT; // initialized in normal way
1190   return Standard_True;
1191 }
1192
1193 // =======================================================================
1194 // function : GetOpenClDeviceInfo
1195 // purpose  : Returns information about device used for computations
1196 // =======================================================================
1197 Standard_Boolean OpenGl_Workspace::GetOpenClDeviceInfo (NCollection_DataMap<TCollection_AsciiString,
1198                                                                             TCollection_AsciiString>& theInfo) const
1199 {
1200   theInfo.Clear();
1201   if (myComputeContext == NULL)
1202   {
1203     return Standard_False;
1204   }
1205
1206   size_t aDevicesSize = 0;
1207   cl_int anError = clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, 0, NULL, &aDevicesSize);
1208   cl_device_id* aDevices = (cl_device_id* )alloca (aDevicesSize);
1209   anError |= clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, aDevicesSize, aDevices, NULL);
1210   if (anError != CL_SUCCESS)
1211   {
1212     return Standard_False;
1213   }
1214
1215   char aDeviceName[256];
1216   anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_NAME, sizeof(aDeviceName), aDeviceName, NULL);
1217   theInfo.Bind ("Name", aDeviceName);
1218
1219   char aDeviceVendor[256];
1220   anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_VENDOR, sizeof(aDeviceVendor), aDeviceVendor, NULL);
1221   theInfo.Bind ("Vendor", aDeviceVendor);
1222
1223   cl_device_type aDeviceType;
1224   anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_TYPE, sizeof(aDeviceType), &aDeviceType, NULL);
1225   theInfo.Bind ("Type", aDeviceType == CL_DEVICE_TYPE_GPU ? "GPU" : "CPU");
1226   return Standard_True;
1227 }
1228
1229 // =======================================================================
1230 // function : ReleaseOpenCL
1231 // purpose  : Releases resources of OpenCL objects
1232 // =======================================================================
1233 void OpenGl_Workspace::ReleaseOpenCL()
1234 {
1235   clReleaseKernel (myRaytraceRenderKernel);
1236   clReleaseKernel (myRaytraceSmoothKernel);
1237
1238   clReleaseProgram (myRaytraceProgram);
1239   clReleaseCommandQueue (myComputeQueue);
1240
1241   clReleaseMemObject (myRaytraceOutputImage);
1242   clReleaseMemObject (myRaytraceEnvironment);
1243   clReleaseMemObject (myRaytraceOutputImageAA);
1244
1245   clReleaseMemObject (myRaytraceMaterialBuffer);
1246   clReleaseMemObject (myRaytraceLightSourceBuffer);
1247
1248   clReleaseMemObject (mySceneNodeInfoBuffer);
1249   clReleaseMemObject (mySceneMinPointBuffer);
1250   clReleaseMemObject (mySceneMaxPointBuffer);
1251
1252   clReleaseMemObject (myObjectNodeInfoBuffer);
1253   clReleaseMemObject (myObjectMinPointBuffer);
1254   clReleaseMemObject (myObjectMaxPointBuffer);
1255
1256   clReleaseMemObject (myGeometryVertexBuffer);
1257   clReleaseMemObject (myGeometryNormalBuffer);
1258   clReleaseMemObject (myGeometryTriangBuffer);
1259
1260   clReleaseContext (myComputeContext);
1261
1262   if (!myGlContext.IsNull())
1263   {
1264     if (!myRaytraceOutputTexture.IsNull())
1265       myGlContext->DelayedRelease (myRaytraceOutputTexture);
1266     myRaytraceOutputTexture.Nullify();
1267
1268     if (!myRaytraceOutputTextureAA.IsNull())
1269       myGlContext->DelayedRelease (myRaytraceOutputTextureAA);
1270     myRaytraceOutputTextureAA.Nullify();
1271   }
1272 }
1273
1274 // =======================================================================
1275 // function : ResizeRaytraceOutputBuffer
1276 // purpose  : Resizes OpenCL output image
1277 // =======================================================================
1278 Standard_Boolean OpenGl_Workspace::ResizeRaytraceOutputBuffer (const cl_int theSizeX,
1279                                                                const cl_int theSizeY)
1280 {
1281   if (myComputeContext == NULL)
1282   {
1283     return Standard_False;
1284   }
1285
1286   if (!myRaytraceOutputTexture.IsNull())
1287   {
1288     Standard_Boolean toResize = myRaytraceOutputTexture->SizeX() != theSizeX ||
1289                                 myRaytraceOutputTexture->SizeY() != theSizeY;
1290
1291     if (!toResize)
1292       return Standard_True;
1293
1294     if (!myGlContext.IsNull())
1295     {
1296       if (!myRaytraceOutputTexture.IsNull())
1297         myGlContext->DelayedRelease (myRaytraceOutputTexture);
1298       if (!myRaytraceOutputTextureAA.IsNull())
1299         myGlContext->DelayedRelease (myRaytraceOutputTextureAA);
1300     }
1301   }
1302
1303   myRaytraceOutputTexture = new OpenGl_Texture();
1304
1305   myRaytraceOutputTexture->Create (myGlContext);
1306   myRaytraceOutputTexture->InitRectangle (myGlContext,
1307     theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>());
1308
1309   myRaytraceOutputTextureAA = new OpenGl_Texture();
1310
1311   myRaytraceOutputTextureAA->Create (myGlContext);
1312   myRaytraceOutputTextureAA->InitRectangle (myGlContext,
1313     theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>());
1314
1315   if (myRaytraceOutputImage != NULL)
1316     clReleaseMemObject (myRaytraceOutputImage);
1317
1318   if (myRaytraceOutputImageAA != NULL)
1319     clReleaseMemObject (myRaytraceOutputImageAA);
1320
1321   cl_int anError = CL_SUCCESS;
1322
1323   myRaytraceOutputImage = clCreateFromGLTexture2D (myComputeContext,
1324     CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTexture->TextureId(), &anError);
1325
1326   if (anError != CL_SUCCESS)
1327   {
1328 #ifdef RAY_TRACE_PRINT_INFO
1329     std::cout << "Error! Failed to create output image!" << std::endl;
1330 #endif
1331     return Standard_False;
1332   }
1333
1334   myRaytraceOutputImageAA = clCreateFromGLTexture2D (myComputeContext,
1335     CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTextureAA->TextureId(), &anError);
1336
1337   if (anError != CL_SUCCESS)
1338   {
1339 #ifdef RAY_TRACE_PRINT_INFO
1340     std::cout << "Error! Failed to create anti-aliased output image!" << std::endl;
1341 #endif
1342     return Standard_False;
1343   }
1344
1345   return Standard_True;
1346 }
1347
1348 // =======================================================================
1349 // function : WriteRaytraceSceneToDevice
1350 // purpose  : Writes scene geometry to OpenCL device
1351 // =======================================================================
1352 Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice()
1353 {
1354   if (myComputeContext == NULL)
1355     return Standard_False;
1356
1357   cl_int anErrorRes = CL_SUCCESS;
1358
1359   if (mySceneNodeInfoBuffer != NULL)
1360     anErrorRes |= clReleaseMemObject (mySceneNodeInfoBuffer);
1361
1362   if (mySceneMinPointBuffer != NULL)
1363     anErrorRes |= clReleaseMemObject (mySceneMinPointBuffer);
1364
1365   if (mySceneMaxPointBuffer != NULL)
1366     anErrorRes |= clReleaseMemObject (mySceneMaxPointBuffer);
1367
1368   if (myObjectNodeInfoBuffer != NULL)
1369     anErrorRes |= clReleaseMemObject (myObjectNodeInfoBuffer);
1370
1371   if (myObjectMinPointBuffer != NULL)
1372     anErrorRes |= clReleaseMemObject (myObjectMinPointBuffer);
1373
1374   if (myObjectMaxPointBuffer != NULL)
1375     anErrorRes |= clReleaseMemObject (myObjectMaxPointBuffer);
1376
1377   if (myGeometryVertexBuffer != NULL)
1378     anErrorRes |= clReleaseMemObject (myGeometryVertexBuffer);
1379
1380   if (myGeometryNormalBuffer != NULL)
1381     anErrorRes |= clReleaseMemObject (myGeometryNormalBuffer);
1382
1383   if (myGeometryTriangBuffer != NULL)
1384     anErrorRes |= clReleaseMemObject (myGeometryTriangBuffer);
1385
1386   if (myRaytraceMaterialBuffer != NULL)
1387     anErrorRes |= clReleaseMemObject (myRaytraceMaterialBuffer);
1388
1389   if (anErrorRes != CL_SUCCESS)
1390   {
1391 #ifdef RAY_TRACE_PRINT_INFO
1392     std::cout << "Error! Failed to release OpenCL buffers" << std::endl;
1393 #endif
1394     return Standard_False;
1395   }
1396
1397   /////////////////////////////////////////////////////////////////////////////
1398   // Create material buffer
1399
1400   const size_t aMaterialBufferSize =
1401     myRaytraceGeometry.Materials.size() != 0 ? myRaytraceGeometry.Materials.size() : 1;
1402
1403   myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext,
1404     CL_MEM_READ_ONLY, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL, &anErrorRes);
1405
1406   if (anErrorRes != CL_SUCCESS)
1407   {
1408 #ifdef RAY_TRACE_PRINT_INFO
1409     std::cout << "Error! Failed to create OpenCL material buffer" << std::endl;
1410 #endif
1411     return Standard_False;
1412   }
1413
1414   /////////////////////////////////////////////////////////////////////////////
1415   // Create BVHs buffers
1416
1417   cl_int anErrorTmp = CL_SUCCESS;
1418
1419   const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = myRaytraceGeometry.BVH();
1420
1421   const size_t aSceneMinPointBufferSize =
1422     aBVH->MinPointBuffer().size() != 0 ? aBVH->MinPointBuffer().size() : 1;
1423
1424   mySceneMinPointBuffer = clCreateBuffer (myComputeContext,
1425     CL_MEM_READ_ONLY, aSceneMinPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp);
1426   anErrorRes |= anErrorTmp;
1427
1428   const size_t aSceneMaxPointBufferSize =
1429     aBVH->MaxPointBuffer().size() != 0 ? aBVH->MaxPointBuffer().size() : 1;
1430
1431   mySceneMaxPointBuffer = clCreateBuffer (myComputeContext,
1432     CL_MEM_READ_ONLY, aSceneMaxPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp);
1433   anErrorRes |= anErrorTmp;
1434
1435   const size_t aSceneNodeInfoBufferSize =
1436     aBVH->NodeInfoBuffer().size() != 0 ? aBVH->NodeInfoBuffer().size() : 1;
1437
1438   mySceneNodeInfoBuffer = clCreateBuffer (myComputeContext,
1439     CL_MEM_READ_ONLY, aSceneNodeInfoBufferSize * sizeof(cl_int4), NULL, &anErrorTmp);
1440   anErrorRes |= anErrorTmp;
1441
1442   if (anErrorRes != CL_SUCCESS)
1443   {
1444 #ifdef RAY_TRACE_PRINT_INFO
1445     std::cout << "Error! Failed to create OpenCL buffers for high-level scene BVH" << std::endl;
1446 #endif
1447     return Standard_False;
1448   }
1449
1450   Standard_Integer aTotalVerticesNb = 0;
1451   Standard_Integer aTotalElementsNb = 0;
1452   Standard_Integer aTotalBVHNodesNb = 0;
1453
1454   for (Standard_Integer anElemIndex = 0; anElemIndex < myRaytraceGeometry.Size(); ++anElemIndex)
1455   {
1456     OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
1457       myRaytraceGeometry.Objects().ChangeValue (anElemIndex).operator->());
1458
1459     Standard_ASSERT_RETURN (aTriangleSet != NULL,
1460       "Error! Failed to get triangulation of OpenGL element", Standard_False);
1461
1462     aTotalVerticesNb += aTriangleSet->Vertices.size();
1463     aTotalElementsNb += aTriangleSet->Elements.size();
1464
1465     Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(),
1466       "Error! Failed to get bottom-level BVH of OpenGL element", Standard_False);
1467
1468     aTotalBVHNodesNb += aTriangleSet->BVH()->NodeInfoBuffer().size();
1469   }
1470
1471   aTotalBVHNodesNb = aTotalBVHNodesNb > 0 ? aTotalBVHNodesNb : 1;
1472
1473   myObjectNodeInfoBuffer = clCreateBuffer (myComputeContext,
1474     CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_int4), NULL, &anErrorTmp);
1475   anErrorRes |= anErrorTmp;
1476
1477   myObjectMinPointBuffer = clCreateBuffer (myComputeContext,
1478     CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp);
1479   anErrorRes |= anErrorTmp;
1480
1481   myObjectMaxPointBuffer = clCreateBuffer (myComputeContext,
1482     CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp);
1483   anErrorRes |= anErrorTmp;
1484
1485   if (anErrorRes != CL_SUCCESS)
1486   {
1487 #ifdef RAY_TRACE_PRINT_INFO
1488     std::cout << "Error! Failed to create OpenCL buffers for bottom-level scene BVHs" << std::endl;
1489 #endif
1490     return Standard_False;
1491   }
1492
1493   /////////////////////////////////////////////////////////////////////////////
1494   // Create geometry buffers
1495
1496   aTotalVerticesNb = aTotalVerticesNb > 0 ? aTotalVerticesNb : 1;
1497
1498   myGeometryVertexBuffer = clCreateBuffer (myComputeContext,
1499     CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp);
1500   anErrorRes |= anErrorTmp;
1501
1502   myGeometryNormalBuffer = clCreateBuffer (myComputeContext,
1503     CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp);
1504   anErrorRes |= anErrorTmp;
1505
1506   aTotalElementsNb = aTotalElementsNb > 0 ? aTotalElementsNb : 1;
1507
1508   myGeometryTriangBuffer = clCreateBuffer (myComputeContext,
1509     CL_MEM_READ_ONLY, aTotalElementsNb * sizeof(cl_int4), NULL, &anErrorTmp);
1510   anErrorRes |= anErrorTmp;
1511
1512   if (anErrorRes != CL_SUCCESS)
1513   {
1514 #ifdef RAY_TRACE_PRINT_INFO
1515     std::cout << "Error! Failed to create OpenCL geometry buffers" << std::endl;
1516 #endif
1517     return Standard_False;
1518   }
1519
1520   /////////////////////////////////////////////////////////////////////////////
1521   // Write BVH and geometry buffers
1522
1523   if (aBVH->NodeInfoBuffer().size() != 0)
1524   {
1525     anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneNodeInfoBuffer, CL_FALSE, 0,
1526       aSceneNodeInfoBufferSize * sizeof(cl_int4), &aBVH->NodeInfoBuffer().front(), 0, NULL, NULL);
1527
1528     anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMinPointBuffer, CL_FALSE, 0,
1529       aSceneMinPointBufferSize * sizeof(cl_float4), &aBVH->MinPointBuffer().front(), 0, NULL, NULL);
1530
1531     anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMaxPointBuffer, CL_FALSE, 0,
1532       aSceneMaxPointBufferSize * sizeof(cl_float4), &aBVH->MaxPointBuffer().front(), 0, NULL, NULL);
1533
1534     anErrorRes |= clFinish (myComputeQueue);
1535
1536     if (anErrorRes != CL_SUCCESS)
1537     {
1538 #ifdef RAY_TRACE_PRINT_INFO
1539       std::cout << "Error! Failed to write OpenCL buffers for high-level scene BVH" << std::endl;
1540 #endif
1541       return Standard_False;
1542     }
1543
1544     for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx)
1545     {
1546       if (!aBVH->IsOuter (aNodeIdx))
1547         continue;
1548
1549       OpenGl_TriangleSet* aTriangleSet = myRaytraceGeometry.TriangleSet (aNodeIdx);
1550
1551       Standard_ASSERT_RETURN (aTriangleSet != NULL,
1552         "Error! Failed to get triangulation of OpenGL element", Standard_False);
1553
1554       const size_t aBVHBuffserSize =
1555         aTriangleSet->BVH()->NodeInfoBuffer().size() != 0 ? aTriangleSet->BVH()->NodeInfoBuffer().size() : 1;
1556
1557       const Standard_Integer aBVHOffset = myRaytraceGeometry.AccelerationOffset (aNodeIdx);
1558
1559       Standard_ASSERT_RETURN (aBVHOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
1560         "Error! Failed to get offset for bottom-level BVH", Standard_False);
1561
1562       anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectNodeInfoBuffer, CL_FALSE, aBVHOffset * sizeof(cl_int4),
1563         aBVHBuffserSize * sizeof(cl_int4), &aTriangleSet->BVH()->NodeInfoBuffer().front(), 0, NULL, NULL);
1564
1565       anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMinPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4),
1566         aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MinPointBuffer().front(), 0, NULL, NULL);
1567
1568       anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMaxPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4),
1569         aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MaxPointBuffer().front(), 0, NULL, NULL);
1570
1571       anErrorRes |= clFinish (myComputeQueue);
1572
1573       if (anErrorRes != CL_SUCCESS)
1574       {
1575 #ifdef RAY_TRACE_PRINT_INFO
1576         std::cout << "Error! Failed to write OpenCL buffers for bottom-level scene BVHs" << std::endl;
1577 #endif
1578         return Standard_False;
1579       }
1580
1581       const Standard_Integer aVerticesOffset = myRaytraceGeometry.VerticesOffset (aNodeIdx);
1582
1583       Standard_ASSERT_RETURN (aVerticesOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
1584         "Error! Failed to get offset for triangulation vertices of OpenGL element", Standard_False);
1585
1586       anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryVertexBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4),
1587         aTriangleSet->Vertices.size() * sizeof(cl_float4), &aTriangleSet->Vertices.front(), 0, NULL, NULL);
1588
1589       anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryNormalBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4),
1590         aTriangleSet->Normals.size() * sizeof(cl_float4), &aTriangleSet->Normals.front(), 0, NULL, NULL);
1591
1592       const Standard_Integer anElementsOffset = myRaytraceGeometry.ElementsOffset (aNodeIdx);
1593
1594       Standard_ASSERT_RETURN (anElementsOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
1595         "Error! Failed to get offset for triangulation elements of OpenGL element", Standard_False);
1596
1597       anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryTriangBuffer, CL_FALSE, anElementsOffset * sizeof(cl_int4),
1598         aTriangleSet->Elements.size() * sizeof(cl_int4), &aTriangleSet->Elements.front(), 0, NULL, NULL);
1599
1600       anErrorRes |= clFinish (myComputeQueue);
1601
1602       if (anErrorRes != CL_SUCCESS)
1603       {
1604 #ifdef RAY_TRACE_PRINT_INFO
1605         std::cout << "Error! Failed to write OpenCL triangulation buffers for OpenGL element" << std::endl;
1606 #endif
1607         return Standard_False;
1608       }
1609     }
1610   }
1611
1612   /////////////////////////////////////////////////////////////////////////////
1613   // Write material buffer
1614
1615   if (myRaytraceGeometry.Materials.size() != 0)
1616   {
1617     const void* aDataPtr = myRaytraceGeometry.Materials.front().Packed();
1618
1619     anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceMaterialBuffer,
1620       CL_FALSE, 0, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr, 0, NULL, NULL);
1621
1622     if (anErrorRes != CL_SUCCESS)
1623     {
1624   #ifdef RAY_TRACE_PRINT_INFO
1625       std::cout << "Error! Failed to write OpenCL material buffer" << std::endl;
1626   #endif
1627       return Standard_False;
1628     }
1629   }
1630
1631   anErrorRes |= clFinish (myComputeQueue);
1632
1633   if (anErrorRes == CL_SUCCESS)
1634   {
1635     myIsRaytraceDataValid = myRaytraceGeometry.Objects().Size() != 0;
1636   }
1637 #ifdef RAY_TRACE_PRINT_INFO
1638   else
1639   {
1640     std::cout << "Error! Failed to set scene data buffers" << std::endl;
1641   }
1642 #endif
1643
1644 #ifdef RAY_TRACE_PRINT_INFO
1645
1646   Standard_ShortReal aMemUsed = 0.f;
1647
1648   for (Standard_Integer anElemIdx = 0; anElemIdx < myRaytraceGeometry.Size(); ++anElemIdx)
1649   {
1650     OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
1651       myRaytraceGeometry.Objects().ChangeValue (anElemIdx).operator->());
1652
1653     aMemUsed += static_cast<Standard_ShortReal> (
1654       aTriangleSet->Vertices.size() * sizeof (BVH_Vec4f));
1655     aMemUsed += static_cast<Standard_ShortReal> (
1656       aTriangleSet->Normals.size() * sizeof (BVH_Vec4f));
1657     aMemUsed += static_cast<Standard_ShortReal> (
1658       aTriangleSet->Elements.size() * sizeof (BVH_Vec4i));
1659
1660     aMemUsed += static_cast<Standard_ShortReal> (
1661       aTriangleSet->BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
1662     aMemUsed += static_cast<Standard_ShortReal> (
1663       aTriangleSet->BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
1664     aMemUsed += static_cast<Standard_ShortReal> (
1665       aTriangleSet->BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
1666   }
1667
1668   aMemUsed += static_cast<Standard_ShortReal> (
1669     myRaytraceGeometry.BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
1670   aMemUsed += static_cast<Standard_ShortReal> (
1671     myRaytraceGeometry.BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
1672   aMemUsed += static_cast<Standard_ShortReal> (
1673     myRaytraceGeometry.BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
1674
1675   std::cout << "GPU Memory Used (MB): ~" << aMemUsed / 1048576 << std::endl;
1676
1677 #endif
1678
1679   return (CL_SUCCESS == anErrorRes);
1680 }
1681
1682 // Use it to estimate the optimal size of OpenCL work group
1683 // #define OPENCL_GROUP_SIZE_TEST
1684
1685 // =======================================================================
1686 // function : RunRaytraceOpenCLKernels
1687 // purpose  : Runs OpenCL ray-tracing kernels
1688 // =======================================================================
1689 Standard_Boolean OpenGl_Workspace::RunRaytraceOpenCLKernels (const Graphic3d_CView&   theCView,
1690                                                              const Standard_ShortReal theOrigins[16],
1691                                                              const Standard_ShortReal theDirects[16],
1692                                                              const Standard_Integer   theSizeX,
1693                                                              const Standard_Integer   theSizeY)
1694 {
1695   if (myRaytraceRenderKernel == NULL || myComputeQueue == NULL)
1696     return Standard_False;
1697
1698   ////////////////////////////////////////////////////////////////////////
1699   // Set kernel arguments
1700
1701   cl_uint anIndex = 0;
1702   cl_int  anError = 0;
1703
1704   cl_int aLightSourceBufferSize = myRaytraceGeometry.Sources.size();
1705
1706   anError |= clSetKernelArg (
1707     myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theSizeX);
1708   anError |= clSetKernelArg (
1709     myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theSizeY);
1710   anError |= clSetKernelArg (
1711     myRaytraceRenderKernel, anIndex++, sizeof(cl_float16), theOrigins);
1712   anError |= clSetKernelArg (
1713     myRaytraceRenderKernel, anIndex++, sizeof(cl_float16), theDirects);
1714   anError |= clSetKernelArg (
1715     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceEnvironment);
1716   anError |= clSetKernelArg (
1717     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImage);
1718   anError |= clSetKernelArg (
1719     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneNodeInfoBuffer);
1720   anError |= clSetKernelArg (
1721     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneMinPointBuffer);
1722   anError |= clSetKernelArg (
1723     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneMaxPointBuffer);
1724   anError |= clSetKernelArg (
1725     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectNodeInfoBuffer);
1726   anError |= clSetKernelArg (
1727     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectMinPointBuffer);
1728   anError |= clSetKernelArg (
1729     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectMaxPointBuffer);
1730   anError |= clSetKernelArg (
1731     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryTriangBuffer);
1732   anError |= clSetKernelArg (
1733     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryVertexBuffer);
1734   anError |= clSetKernelArg (
1735     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryNormalBuffer);
1736   anError |= clSetKernelArg (
1737     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceLightSourceBuffer);
1738   anError |= clSetKernelArg (
1739     myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceMaterialBuffer);
1740   anError |= clSetKernelArg (
1741     myRaytraceRenderKernel, anIndex++, sizeof(cl_float4), &myRaytraceGeometry.GlobalAmbient);
1742   anError |= clSetKernelArg (
1743     myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &aLightSourceBufferSize);
1744   anError |= clSetKernelArg (
1745     myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theCView.IsShadowsEnabled);
1746   anError |= clSetKernelArg (
1747     myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theCView.IsReflectionsEnabled);
1748   anError |= clSetKernelArg (
1749     myRaytraceRenderKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneEpsilon);
1750   anError |= clSetKernelArg (
1751     myRaytraceRenderKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneRadius);
1752
1753   if (anError != CL_SUCCESS)
1754   {
1755     const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of ray-tracing kernel!";
1756
1757     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1758       GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
1759
1760     return Standard_False;
1761   }
1762
1763   // Second-pass 'smoothing' kernel runs only if anti-aliasing is enabled
1764   if (theCView.IsAntialiasingEnabled)
1765   {
1766     anIndex = 0;
1767
1768     anError |= clSetKernelArg (
1769       myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theSizeX);
1770     anError |= clSetKernelArg (
1771       myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theSizeY);
1772     anError |= clSetKernelArg (
1773       myRaytraceSmoothKernel, anIndex++, sizeof(cl_float16), theOrigins);
1774     anError |= clSetKernelArg (
1775       myRaytraceSmoothKernel, anIndex++, sizeof(cl_float16), theDirects);
1776     anError |= clSetKernelArg (
1777       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImage);
1778     anError |= clSetKernelArg (
1779       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceEnvironment);
1780     anError |= clSetKernelArg (
1781       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImageAA);
1782     anError |= clSetKernelArg (
1783       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneNodeInfoBuffer);
1784     anError |= clSetKernelArg (
1785       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneMinPointBuffer);
1786     anError |= clSetKernelArg (
1787       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneMaxPointBuffer);
1788     anError |= clSetKernelArg (
1789       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectNodeInfoBuffer);
1790     anError |= clSetKernelArg (
1791       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectMinPointBuffer);
1792     anError |= clSetKernelArg (
1793       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectMaxPointBuffer);
1794     anError |= clSetKernelArg (
1795       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryTriangBuffer);
1796     anError |= clSetKernelArg (
1797       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryVertexBuffer);
1798     anError |= clSetKernelArg (
1799       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryNormalBuffer);
1800     anError |= clSetKernelArg (
1801       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceLightSourceBuffer);
1802     anError |= clSetKernelArg (
1803       myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceMaterialBuffer);
1804     anError |= clSetKernelArg (
1805       myRaytraceSmoothKernel, anIndex++, sizeof(cl_float4), &myRaytraceGeometry.GlobalAmbient);
1806     anError |= clSetKernelArg (
1807       myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &aLightSourceBufferSize);
1808     anError |= clSetKernelArg (
1809       myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theCView.IsShadowsEnabled);
1810     anError |= clSetKernelArg (
1811       myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theCView.IsReflectionsEnabled);
1812     anError |= clSetKernelArg (
1813       myRaytraceSmoothKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneEpsilon);
1814     anError |= clSetKernelArg (
1815       myRaytraceSmoothKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneRadius);
1816
1817     if (anError != CL_SUCCESS)
1818     {
1819       const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of smoothing kernel!";
1820
1821       myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1822         GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
1823
1824       return Standard_False;
1825     }
1826   }
1827
1828   ////////////////////////////////////////////////////////////////////////
1829   // Set work size
1830
1831   size_t aLocWorkSize[] = { myIsAmdComputePlatform ? 2 : 4, 32 };
1832
1833 #ifdef OPENCL_GROUP_SIZE_TEST
1834   for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1)
1835   for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1)
1836   {
1837     aLocWorkSize[0] = aLocX;
1838     aLocWorkSize[1] = aLocY;
1839 #endif
1840
1841     size_t aWorkSizeX = theSizeX;
1842     if (aWorkSizeX % aLocWorkSize[0] != 0)
1843       aWorkSizeX += aLocWorkSize[0] - aWorkSizeX % aLocWorkSize[0];
1844
1845     size_t aWokrSizeY = theSizeY;
1846     if (aWokrSizeY % aLocWorkSize[1] != 0 )
1847       aWokrSizeY += aLocWorkSize[1] - aWokrSizeY % aLocWorkSize[1];
1848
1849     size_t aTotWorkSize[] = { aWorkSizeX, aWokrSizeY };
1850
1851     cl_event anEvent = NULL, anEventSmooth = NULL;
1852
1853     anError = clEnqueueNDRangeKernel (myComputeQueue,
1854       myRaytraceRenderKernel, 2, NULL, aTotWorkSize, aLocWorkSize, 0, NULL, &anEvent);
1855
1856     if (anError != CL_SUCCESS)
1857     {
1858       const TCollection_ExtendedString aMessage = "Error! Failed to execute the ray-tracing kernel!";
1859
1860       myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1861         GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
1862
1863       return Standard_False;
1864     }
1865
1866     clWaitForEvents (1, &anEvent);
1867
1868     if (theCView.IsAntialiasingEnabled)
1869     {
1870       size_t aLocWorkSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4,
1871                                       myIsAmdComputePlatform ? 8 : 32 };
1872
1873 #ifdef OPENCL_GROUP_SIZE_TEST
1874       aLocWorkSizeSmooth[0] = aLocX;
1875       aLocWorkSizeSmooth[1] = aLocY;
1876 #endif
1877
1878       aWorkSizeX = theSizeX;
1879       if (aWorkSizeX % aLocWorkSizeSmooth[0] != 0)
1880         aWorkSizeX += aLocWorkSizeSmooth[0] - aWorkSizeX % aLocWorkSizeSmooth[0];
1881
1882       size_t aWokrSizeY = theSizeY;
1883       if (aWokrSizeY % aLocWorkSizeSmooth[1] != 0 )
1884         aWokrSizeY += aLocWorkSizeSmooth[1] - aWokrSizeY % aLocWorkSizeSmooth[1];
1885
1886       size_t aTotWorkSizeSmooth [] = { aWorkSizeX, aWokrSizeY };
1887
1888       anError = clEnqueueNDRangeKernel (myComputeQueue, myRaytraceSmoothKernel,
1889         2, NULL, aTotWorkSizeSmooth, aLocWorkSizeSmooth, 0, NULL, &anEventSmooth);
1890
1891       clWaitForEvents (1, &anEventSmooth);
1892
1893       if (anError != CL_SUCCESS)
1894       {
1895         const TCollection_ExtendedString aMessage = "Error! Failed to execute the smoothing kernel";
1896
1897         myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1898           GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
1899
1900         return Standard_False;
1901       }
1902     }
1903
1904 #if defined (RAY_TRACE_PRINT_INFO) || defined (OPENCL_GROUP_SIZE_TEST)
1905
1906     static cl_ulong ttt1 = 10000000000;
1907     static cl_ulong ttt2 = 10000000000;
1908
1909     cl_ulong aBegTime = 0;
1910     cl_ulong aEndTime = 0;
1911
1912     clGetEventProfilingInfo (anEvent,
1913       CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL);
1914     clGetEventProfilingInfo (anEvent,
1915       CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL);
1916
1917     ttt1 = aEndTime - aBegTime < ttt1 ? aEndTime - aBegTime : ttt1;
1918
1919     std::cout << "\tRender time (ms): " << ttt1 / 1e6f << std::endl;
1920
1921     if (theCView.IsAntialiasingEnabled)
1922     {
1923       clGetEventProfilingInfo (anEventSmooth,
1924         CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL);
1925       clGetEventProfilingInfo (anEventSmooth,
1926         CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL);
1927
1928       ttt2 = aEndTime - aBegTime < ttt2 ? aEndTime - aBegTime : ttt2;
1929
1930       std::cout << "\tSmooth time (ms): " << ttt2 / 1e6f << std::endl;
1931     }
1932
1933 #endif
1934
1935     if (anEvent != NULL)
1936       clReleaseEvent (anEvent);
1937
1938     if (anEventSmooth != NULL)
1939       clReleaseEvent (anEventSmooth);
1940
1941 #ifdef OPENCL_GROUP_SIZE_TEST
1942   }
1943 #endif
1944
1945   return Standard_True;
1946 }
1947
1948 // =======================================================================
1949 // function : ComputeInverseMatrix
1950 // purpose  : Computes inversion of 4x4 floating-point matrix
1951 // =======================================================================
1952 template <typename T>
1953 void ComputeInverseMatrix (const T m[16], T inv[16])
1954 {
1955   inv[ 0] = m[ 5] * (m[10] * m[15] - m[11] * m[14]) -
1956             m[ 9] * (m[ 6] * m[15] - m[ 7] * m[14]) -
1957             m[13] * (m[ 7] * m[10] - m[ 6] * m[11]);
1958
1959   inv[ 1] = m[ 1] * (m[11] * m[14] - m[10] * m[15]) -
1960             m[ 9] * (m[ 3] * m[14] - m[ 2] * m[15]) -
1961             m[13] * (m[ 2] * m[11] - m[ 3] * m[10]);
1962
1963   inv[ 2] = m[ 1] * (m[ 6] * m[15] - m[ 7] * m[14]) -
1964             m[ 5] * (m[ 2] * m[15] - m[ 3] * m[14]) -
1965             m[13] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]);
1966
1967   inv[ 3] = m[ 1] * (m[ 7] * m[10] - m[ 6] * m[11]) -
1968             m[ 5] * (m[ 3] * m[10] - m[ 2] * m[11]) -
1969             m[ 9] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]);
1970
1971   inv[ 4] = m[ 4] * (m[11] * m[14] - m[10] * m[15]) -
1972             m[ 8] * (m[ 7] * m[14] - m[ 6] * m[15]) -
1973             m[12] * (m[ 6] * m[11] - m[ 7] * m[10]);
1974
1975   inv[ 5] = m[ 0] * (m[10] * m[15] - m[11] * m[14]) -
1976             m[ 8] * (m[ 2] * m[15] - m[ 3] * m[14]) -
1977             m[12] * (m[ 3] * m[10] - m[ 2] * m[11]);
1978
1979   inv[ 6] = m[ 0] * (m[ 7] * m[14] - m[ 6] * m[15]) -
1980             m[ 4] * (m[ 3] * m[14] - m[ 2] * m[15]) -
1981             m[12] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]);
1982
1983   inv[ 7] = m[ 0] * (m[ 6] * m[11] - m[ 7] * m[10]) -
1984             m[ 4] * (m[ 2] * m[11] - m[ 3] * m[10]) -
1985             m[ 8] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]);
1986
1987   inv[ 8] = m[ 4] * (m[ 9] * m[15] - m[11] * m[13]) -
1988             m[ 8] * (m[ 5] * m[15] - m[ 7] * m[13]) -
1989             m[12] * (m[ 7] * m[ 9] - m[ 5] * m[11]);
1990
1991   inv[ 9] = m[ 0] * (m[11] * m[13] - m[ 9] * m[15]) -
1992             m[ 8] * (m[ 3] * m[13] - m[ 1] * m[15]) -
1993             m[12] * (m[ 1] * m[11] - m[ 3] * m[ 9]);
1994
1995   inv[10] = m[ 0] * (m[ 5] * m[15] - m[ 7] * m[13]) -
1996             m[ 4] * (m[ 1] * m[15] - m[ 3] * m[13]) -
1997             m[12] * (m[ 3] * m[ 5] - m[ 1] * m[ 7]);
1998
1999   inv[11] = m[ 0] * (m[ 7] * m[ 9] - m[ 5] * m[11]) -
2000             m[ 4] * (m[ 3] * m[ 9] - m[ 1] * m[11]) -
2001             m[ 8] * (m[ 1] * m[ 7] - m[ 3] * m[ 5]);
2002
2003   inv[12] = m[ 4] * (m[10] * m[13] - m[ 9] * m[14]) -
2004             m[ 8] * (m[ 6] * m[13] - m[ 5] * m[14]) -
2005             m[12] * (m[ 5] * m[10] - m[ 6] * m[ 9]);
2006
2007   inv[13] = m[ 0] * (m[ 9] * m[14] - m[10] * m[13]) -
2008             m[ 8] * (m[ 1] * m[14] - m[ 2] * m[13]) -
2009             m[12] * (m[ 2] * m[ 9] - m[ 1] * m[10]);
2010
2011   inv[14] = m[ 0] * (m[ 6] * m[13] - m[ 5] * m[14]) -
2012             m[ 4] * (m[ 2] * m[13] - m[ 1] * m[14]) -
2013             m[12] * (m[ 1] * m[ 6] - m[ 2] * m[ 5]);
2014
2015   inv[15] = m[ 0] * (m[ 5] * m[10] - m[ 6] * m[ 9]) -
2016             m[ 4] * (m[ 1] * m[10] - m[ 2] * m[ 9]) -
2017             m[ 8] * (m[ 2] * m[ 5] - m[ 1] * m[ 6]);
2018
2019   T det = m[0] * inv[ 0] +
2020           m[1] * inv[ 4] +
2021           m[2] * inv[ 8] +
2022           m[3] * inv[12];
2023
2024   if (det == T (0.0)) return;
2025
2026   det = T (1.0) / det;
2027
2028   for (Standard_Integer i = 0; i < 16; ++i)
2029     inv[i] *= det;
2030 }
2031
2032 // =======================================================================
2033 // function : GenerateCornerRays
2034 // purpose  : Generates primary rays for corners of screen quad
2035 // =======================================================================
2036 void GenerateCornerRays (const GLdouble theInvModelProj[16],
2037                          cl_float       theOrigins[16],
2038                          cl_float       theDirects[16])
2039 {
2040   Standard_Integer aOriginIndex = 0;
2041   Standard_Integer aDirectIndex = 0;
2042
2043   for (Standard_Integer y = -1; y <= 1; y += 2)
2044   {
2045     for (Standard_Integer x = -1; x <= 1; x += 2)
2046     {
2047       BVH_Vec4f aOrigin (float(x),
2048                          float(y),
2049                          -1.f,
2050                          1.f);
2051
2052       aOrigin = MatVecMult (theInvModelProj, aOrigin);
2053       aOrigin.x() = aOrigin.x() / aOrigin.w();
2054       aOrigin.y() = aOrigin.y() / aOrigin.w();
2055       aOrigin.z() = aOrigin.z() / aOrigin.w();
2056       aOrigin.w() = 1.f;
2057
2058       BVH_Vec4f aDirect (float(x),
2059                          float(y),
2060                          1.f,
2061                          1.f);
2062
2063       aDirect = MatVecMult (theInvModelProj, aDirect);
2064       aDirect.x() = aDirect.x() / aDirect.w();
2065       aDirect.y() = aDirect.y() / aDirect.w();
2066       aDirect.z() = aDirect.z() / aDirect.w();
2067       aDirect.w() = 1.f;
2068
2069       aDirect = aDirect - aOrigin;
2070
2071       GLdouble aInvLen = 1.f / sqrt (aDirect.x() * aDirect.x() +
2072                                      aDirect.y() * aDirect.y() +
2073                                      aDirect.z() * aDirect.z());
2074
2075       theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.x());
2076       theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.y());
2077       theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.z());
2078       theOrigins [aOriginIndex++] = 1.f;
2079
2080       theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.x() * aInvLen);
2081       theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.y() * aInvLen);
2082       theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.z() * aInvLen);
2083       theDirects [aDirectIndex++] = 0.f;
2084     }
2085   }
2086 }
2087
2088 // =======================================================================
2089 // function : Raytrace
2090 // purpose  : Redraws the window using OpenCL ray tracing
2091 // =======================================================================
2092 Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView,
2093                                              const Standard_Integer theSizeX,
2094                                              const Standard_Integer theSizeY,
2095                                              const Tint             theToSwap)
2096 {
2097   if (!InitOpenCL())
2098     return Standard_False;
2099
2100   if (!ResizeRaytraceOutputBuffer (theSizeX, theSizeY))
2101     return Standard_False;
2102
2103   if (!UpdateRaytraceEnvironmentMap())
2104     return Standard_False;
2105
2106   if (!UpdateRaytraceGeometry (Standard_True))
2107     return Standard_False;
2108
2109   // Get model-view and projection matrices
2110   TColStd_Array2OfReal theOrientation (0, 3, 0, 3);
2111   TColStd_Array2OfReal theViewMapping (0, 3, 0, 3);
2112
2113   myView->GetMatrices (theOrientation, theViewMapping);
2114
2115   GLdouble aOrientationMatrix[16];
2116   GLdouble aViewMappingMatrix[16];
2117   GLdouble aOrientationInvers[16];
2118
2119   for (Standard_Integer j = 0; j < 4; ++j)
2120     for (Standard_Integer i = 0; i < 4; ++i)
2121     {
2122       aOrientationMatrix [4 * j + i] = theOrientation (i, j);
2123       aViewMappingMatrix [4 * j + i] = theViewMapping (i, j);
2124     }
2125
2126   ComputeInverseMatrix (aOrientationMatrix, aOrientationInvers);
2127
2128   if (!UpdateRaytraceLightSources (aOrientationInvers))
2129     return Standard_False;
2130
2131   // Generate primary rays for corners of the screen quad
2132   glMatrixMode (GL_MODELVIEW);
2133
2134   glLoadMatrixd (aViewMappingMatrix);
2135   glMultMatrixd (aOrientationMatrix);
2136
2137   GLdouble aModelProject[16];
2138   GLdouble aInvModelProj[16];
2139
2140   glGetDoublev (GL_MODELVIEW_MATRIX, aModelProject);
2141
2142   ComputeInverseMatrix (aModelProject, aInvModelProj);
2143
2144   GLfloat aOrigins[16];
2145   GLfloat aDirects[16];
2146
2147   GenerateCornerRays (aInvModelProj,
2148                       aOrigins,
2149                       aDirects);
2150
2151   // Compute ray-traced image using OpenCL kernel
2152   cl_mem anImages[] = { myRaytraceOutputImage, myRaytraceOutputImageAA };
2153   cl_int anError = clEnqueueAcquireGLObjects (myComputeQueue,
2154                                               2, anImages,
2155                                               0, NULL, NULL);
2156   clFinish (myComputeQueue);
2157
2158   if (myIsRaytraceDataValid)
2159   {
2160     RunRaytraceOpenCLKernels (theCView,
2161                               aOrigins,
2162                               aDirects,
2163                               theSizeX,
2164                               theSizeY);
2165   }
2166
2167   anError |= clEnqueueReleaseGLObjects (myComputeQueue,
2168                                         2, anImages,
2169                                         0, NULL, NULL);
2170   clFinish (myComputeQueue);
2171
2172   // Draw background
2173   glPushAttrib (GL_ENABLE_BIT |
2174                 GL_CURRENT_BIT |
2175                 GL_COLOR_BUFFER_BIT |
2176                 GL_DEPTH_BUFFER_BIT);
2177
2178   glDisable (GL_DEPTH_TEST);
2179
2180   if (NamedStatus & OPENGL_NS_WHITEBACK)
2181   {
2182     glClearColor (1.0f, 1.0f, 1.0f, 1.0f);
2183   }
2184   else
2185   {
2186     glClearColor (myBgColor.rgb[0],
2187                   myBgColor.rgb[1],
2188                   myBgColor.rgb[2],
2189                   1.0f);
2190   }
2191
2192   glClear (GL_COLOR_BUFFER_BIT);
2193
2194   Handle(OpenGl_Workspace) aWorkspace (this);
2195   myView->DrawBackground (aWorkspace);
2196
2197   // Draw dummy quad to show result image
2198   glEnable (GL_COLOR_MATERIAL);
2199   glEnable (GL_BLEND);
2200
2201   glDisable (GL_DEPTH_TEST);
2202
2203   glBlendFunc (GL_ONE, GL_SRC_ALPHA);
2204
2205   glEnable (GL_TEXTURE_RECTANGLE);
2206
2207   glMatrixMode (GL_PROJECTION);
2208   glLoadIdentity();
2209
2210   glMatrixMode (GL_MODELVIEW);
2211   glLoadIdentity();
2212
2213   glColor3f (1.0f, 1.0f, 1.0f);
2214
2215   if (!theCView.IsAntialiasingEnabled)
2216     myRaytraceOutputTexture->Bind (myGlContext);
2217   else
2218     myRaytraceOutputTextureAA->Bind (myGlContext);
2219
2220   if (myIsRaytraceDataValid)
2221   {
2222     glBegin (GL_QUADS);
2223     {
2224       glTexCoord2i (       0,        0);   glVertex2f (-1.f, -1.f);
2225       glTexCoord2i (       0, theSizeY);   glVertex2f (-1.f,  1.f);
2226       glTexCoord2i (theSizeX, theSizeY);   glVertex2f ( 1.f,  1.f);
2227       glTexCoord2i (theSizeX,        0);   glVertex2f ( 1.f, -1.f);
2228     }
2229     glEnd();
2230   }
2231
2232   glPopAttrib();
2233
2234   // Swap the buffers
2235   if (theToSwap)
2236   {
2237     GetGlContext()->SwapBuffers();
2238     myBackBufferRestored = Standard_False;
2239   }
2240   else
2241     glFlush();
2242
2243   return Standard_True;
2244 }
2245
2246 #endif