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