src/OpenGl/OpenGl_Workspace_Raytrace.cxx

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