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