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   Standard_Integer aSizeX = 1;
  91   Standard_Integer 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 (Standard_Integer aPixel = 0; aPixel < aSizeX * aSizeY * 4; ++aPixel)
 128       aPixelData[aPixel] = 0.f;
 129   }
 130
 131   size_t anImageOffset[] = { 0,
 132                              0,
 133                              0 };
 134
 135   size_t anImageRegion[] = { aSizeX,
 136                              aSizeY,
 137                              1 };
 138
 139   anError |= clEnqueueWriteImage (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 (Standard_Integer anIndex = 0; anIndex < aStructArray.Length(); ++anIndex)
 199     {
 200       OpenGl_SequenceOfStructure::Iterator aStructIt;
 201
 202       for (aStructIt.Init (aStructArray (anIndex)); aStructIt.More(); aStructIt.Next())
 203       {
 204         const OpenGl_Structure* aStructure = aStructIt.Value();
 205
 206         if (theCheck)
 207         {
 208           if (CheckRaytraceStructure (aStructure))
 209           {
 210             return UpdateRaytraceGeometry (Standard_False);
 211           }
 212         }
 213         else
 214         {
 215           if (!aStructure->IsRaytracable())
 216             continue;
 217
 218           if (aStructure->Transformation()->mat != NULL)
 219           {
 220             if (aTransform == NULL)
 221               aTransform = new float[16];
 222
 223             for (Standard_Integer i = 0; i < 4; ++i)
 224               for (Standard_Integer j = 0; j < 4; ++j)
 225               {
 226                 aTransform[j * 4 + i] = aStructure->Transformation()->mat[i][j];
 227               }
 228           }
 229
 230           AddRaytraceStructure (aStructure, aTransform, anElements);
 231         }
 232       }
 233     }
 234   }
 235
 236   if (!theCheck)
 237   {
 238     // Actualize the hash map of structures -- remove out-of-date records
 239     std::map<const OpenGl_Structure*, Standard_Size>::iterator anIter = myStructureStates.begin();
 240
 241     while (anIter != myStructureStates.end())
 242     {
 243       if (anElements.find (anIter->first) == anElements.end())
 244       {
 245         myStructureStates.erase (anIter++);
 246       }
 247       else
 248       {
 249         ++anIter;
 250       }
 251     }
 252
 253     // Actualize OpenGL layer list state
 254     myLayersModificationStatus = myView->LayerList().ModificationState();
 255
 256
 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   Standard_Integer aStructMatID = -1;
 389
 390   if (theStructure->AspectFace() != NULL)
 391   {
 392     aStructMatID = static_cast<Standard_Integer> (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     Standard_Integer aGroupMatID = -1;
 406
 407     if (anItg.Value()->AspectFace() != NULL)
 408     {
 409       aGroupMatID = static_cast<Standard_Integer> (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     Standard_Integer aMatID = aGroupMatID < 0 ? aStructMatID : aGroupMatID;
 418
 419     if (aStructMatID < 0 && aGroupMatID < 0)
 420     {
 421       aMatID = static_cast<Standard_Integer> (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<Standard_Integer> (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 (Standard_Integer i = 0; i < 4; ++i)
 469         for (Standard_Integer 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                                                               Standard_Integer       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 Standard_Integer aFirstVert = static_cast<Standard_Integer> (myRaytraceSceneData.Vertices.size());
 534
 535   for (Standard_Integer 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 (Standard_Integer 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     Standard_Integer aVertOffset = 0;
 581
 582     for (Standard_Integer aBound = 0; aBound < theArray->num_bounds; ++aBound)
 583     {
 584       const Standard_Integer 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 Standard_Integer 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                                                              Standard_Integer       theFirstVert,
 619                                                              Standard_Integer       theVertOffset,
 620                                                              Standard_Integer       theVertNum,
 621                                                              Standard_Integer       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                                                              Standard_Integer       theFirstVert,
 642                                                              Standard_Integer       theVertOffset,
 643                                                              Standard_Integer       theVertNum,
 644                                                              Standard_Integer       theMatID)
 645 {
 646   if (theVertNum < 3)
 647     return Standard_True;
 648
 649   if (theArray->num_edges > 0)
 650   {
 651     for (Standard_Integer 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 (Standard_Integer 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                                                                 Standard_Integer       theFirstVert,
 679                                                                 Standard_Integer       theVertOffset,
 680                                                                 Standard_Integer       theVertNum,
 681                                                                 Standard_Integer       theMatID)
 682 {
 683   if (theVertNum < 3)
 684     return Standard_True;
 685
 686   if (theArray->num_edges > 0)
 687   {
 688     for (Standard_Integer 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 (Standard_Integer 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                                                                   Standard_Integer       theFirstVert,
 716                                                                   Standard_Integer       theVertOffset,
 717                                                                   Standard_Integer       theVertNum,
 718                                                                   Standard_Integer       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 (Standard_Integer 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 (Standard_Integer 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                                                                Standard_Integer       theFirstVert,
 765                                                                Standard_Integer       theVertOffset,
 766                                                                Standard_Integer       theVertNum,
 767                                                                Standard_Integer       theMatID)
 768 {
 769   if (theVertNum < 4)
 770     return Standard_True;
 771
 772   if (theArray->num_edges > 0)
 773   {
 774     for (Standard_Integer 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 (Standard_Integer 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                                                                     Standard_Integer       theFirstVert,
 812                                                                     Standard_Integer       theVertOffset,
 813                                                                     Standard_Integer       theVertNum,
 814                                                                     Standard_Integer       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 (Standard_Integer 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 (Standard_Integer 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                                                             Standard_Integer       theFirstVert,
 883                                                             Standard_Integer       theVertOffset,
 884                                                             Standard_Integer       theVertNum,
 885                                                             Standard_Integer       theMatID)
 886 {
 887   if (theArray->num_vertexs < 3)
 888     return Standard_True;
 889
 890   if (theArray->edges != NULL)
 891   {
 892     for (Standard_Integer 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 (Standard_Integer 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 (myRaytraceOutputImageAA);
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 (!myGlContext.IsNull())
1317   {
1318     if (!myRaytraceOutputTexture.IsNull())
1319       myGlContext->DelayedRelease (myRaytraceOutputTexture);
1320     myRaytraceOutputTexture.Nullify();
1321
1322     if (!myRaytraceOutputTextureAA.IsNull())
1323       myGlContext->DelayedRelease (myRaytraceOutputTextureAA);
1324     myRaytraceOutputTextureAA.Nullify();
1325   }
1326 }
1327
1328 // =======================================================================
1329 // function : ResizeRaytraceOutputBuffer
1330 // purpose  : Resizes OpenCL output image
1331 // =======================================================================
1332 Standard_Boolean OpenGl_Workspace::ResizeRaytraceOutputBuffer (const cl_int theSizeX,
1333                                                                const cl_int theSizeY)
1334 {
1335   if (myComputeContext == NULL)
1336   {
1337     return Standard_False;
1338   }
1339
1340   if (!myRaytraceOutputTexture.IsNull())
1341   {
1342     Standard_Boolean toResize = myRaytraceOutputTexture->SizeX() != theSizeX ||
1343                                 myRaytraceOutputTexture->SizeY() != theSizeY;
1344
1345     if (!toResize)
1346       return Standard_True;
1347
1348     if (!myGlContext.IsNull())
1349     {
1350       if (!myRaytraceOutputTexture.IsNull())
1351         myGlContext->DelayedRelease (myRaytraceOutputTexture);
1352       if (!myRaytraceOutputTextureAA.IsNull())
1353         myGlContext->DelayedRelease (myRaytraceOutputTextureAA);
1354     }
1355   }
1356
1357   myRaytraceOutputTexture = new OpenGl_Texture();
1358
1359   myRaytraceOutputTexture->Create (myGlContext);
1360   myRaytraceOutputTexture->InitRectangle (myGlContext,
1361     theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>());
1362
1363   myRaytraceOutputTextureAA = new OpenGl_Texture();
1364
1365   myRaytraceOutputTextureAA->Create (myGlContext);
1366   myRaytraceOutputTextureAA->InitRectangle (myGlContext,
1367     theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>());
1368
1369   if (myRaytraceOutputImage != NULL)
1370     clReleaseMemObject (myRaytraceOutputImage);
1371
1372   if (myRaytraceOutputImageAA != NULL)
1373     clReleaseMemObject (myRaytraceOutputImageAA);
1374
1375   cl_int anError = CL_SUCCESS;
1376
1377   myRaytraceOutputImage = clCreateFromGLTexture2D (myComputeContext,
1378     CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTexture->TextureId(), &anError);
1379
1380   if (anError != CL_SUCCESS)
1381   {
1382 #ifdef RAY_TRACE_PRINT_INFO
1383     std::cout << "Error! Failed to create output image!" << std::endl;
1384 #endif
1385     return Standard_False;
1386   }
1387
1388   myRaytraceOutputImageAA = clCreateFromGLTexture2D (myComputeContext,
1389     CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTextureAA->TextureId(), &anError);
1390
1391   if (anError != CL_SUCCESS)
1392   {
1393 #ifdef RAY_TRACE_PRINT_INFO
1394     std::cout << "Error! Failed to create anti-aliased output image!" << std::endl;
1395 #endif
1396     return Standard_False;
1397   }
1398
1399   return Standard_True;
1400 }
1401
1402 // =======================================================================
1403 // function : WriteRaytraceSceneToDevice
1404 // purpose  : Writes scene geometry to OpenCl device
1405 // =======================================================================
1406 Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice()
1407 {
1408   if (myComputeContext == NULL)
1409     return Standard_False;
1410
1411   cl_int anError = CL_SUCCESS;
1412
1413   if (myRaytraceNormalBuffer != NULL)
1414     anError |= clReleaseMemObject (myRaytraceNormalBuffer);
1415
1416   if (myRaytraceVertexBuffer != NULL)
1417     anError |= clReleaseMemObject (myRaytraceVertexBuffer);
1418
1419   if (myRaytraceTriangleBuffer != NULL)
1420     anError |= clReleaseMemObject (myRaytraceTriangleBuffer);
1421
1422   if (myRaytraceNodeMinPointBuffer != NULL)
1423     anError |= clReleaseMemObject (myRaytraceNodeMinPointBuffer);
1424
1425   if (myRaytraceNodeMaxPointBuffer != NULL)
1426     anError |= clReleaseMemObject (myRaytraceNodeMaxPointBuffer);
1427
1428   if (myRaytraceNodeDataRcrdBuffer != NULL)
1429     anError |= clReleaseMemObject (myRaytraceNodeDataRcrdBuffer);
1430
1431   if (myRaytraceMaterialBuffer != NULL)
1432     anError |= clReleaseMemObject (myRaytraceMaterialBuffer);
1433
1434   if (anError != CL_SUCCESS)
1435   {
1436 #ifdef RAY_TRACE_PRINT_INFO
1437     std::cout << "Error! Failed to release OpenCL scene buffers!" << std::endl;
1438 #endif
1439     return Standard_False;
1440   }
1441
1442   // Create geometry buffers
1443   cl_int anErrorTemp = CL_SUCCESS;
1444   const size_t myVertexBufferSize = myRaytraceSceneData.Vertices.size() > 0
1445                                   ? myRaytraceSceneData.Vertices.size() : 1;
1446
1447   myRaytraceVertexBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
1448                                            myVertexBufferSize * sizeof(cl_float4), NULL, &anErrorTemp);
1449   anError |= anErrorTemp;
1450
1451   const size_t myNormalBufferSize = myRaytraceSceneData.Normals.size() > 0
1452                                   ? myRaytraceSceneData.Normals.size() : 1;
1453   myRaytraceNormalBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
1454                                            myNormalBufferSize * sizeof(cl_float4), NULL, &anErrorTemp);
1455   anError |= anErrorTemp;
1456
1457   const size_t myTriangleBufferSize = myRaytraceSceneData.Triangles.size() > 0
1458                                     ? myRaytraceSceneData.Triangles.size() : 1;
1459   myRaytraceTriangleBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
1460                                              myTriangleBufferSize * sizeof(cl_int4), NULL, &anErrorTemp);
1461   anError |= anErrorTemp;
1462   if (anError != CL_SUCCESS)
1463   {
1464 #ifdef RAY_TRACE_PRINT_INFO
1465     std::cout << "Error! Failed to create OpenCL geometry buffers!" << std::endl;
1466 #endif
1467     return Standard_False;
1468   }
1469
1470   // Create material buffer
1471   const size_t myMaterialBufferSize = myRaytraceSceneData.Materials.size() > 0
1472                                     ? myRaytraceSceneData.Materials.size() : 1;
1473   myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
1474                                              myMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL,
1475                                              &anErrorTemp);
1476   if (anErrorTemp != CL_SUCCESS)
1477   {
1478 #ifdef RAY_TRACE_PRINT_INFO
1479     std::cout << "Error! Failed to create OpenCL material buffer!" << std::endl;
1480 #endif
1481     return Standard_False;
1482   }
1483
1484   // Create BVH buffers
1485   OpenGl_BVH aTree = myBVHBuilder.Tree();
1486   const size_t myNodeMinPointBufferSize = aTree.MinPointBuffer().size() > 0
1487                                         ? aTree.MinPointBuffer().size() : 1;
1488   myRaytraceNodeMinPointBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
1489                                                  myNodeMinPointBufferSize * sizeof(cl_float4), NULL,
1490                                                  &anErrorTemp);
1491   anError |= anErrorTemp;
1492
1493   const size_t myNodeMaxPointBufferSize = aTree.MaxPointBuffer().size() > 0
1494                                         ? aTree.MaxPointBuffer().size() : 1;
1495   myRaytraceNodeMaxPointBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
1496                                                  myNodeMaxPointBufferSize * sizeof(cl_float4), NULL,
1497                                                  &anError);
1498   anError |= anErrorTemp;
1499
1500   const size_t myNodeDataRecordBufferSize = aTree.DataRcrdBuffer().size() > 0
1501                                           ? aTree.DataRcrdBuffer().size() : 1;
1502   myRaytraceNodeDataRcrdBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
1503                                                  myNodeDataRecordBufferSize * sizeof(cl_int4), NULL,
1504                                                  &anError);
1505   anError |= anErrorTemp;
1506   if (anError != CL_SUCCESS)
1507   {
1508 #ifdef RAY_TRACE_PRINT_INFO
1509     std::cout << "Error! Failed to create OpenCL BVH buffers!" << std::endl;
1510 #endif
1511     return Standard_False;
1512   }
1513
1514   // Write scene geometry buffers
1515   if (myRaytraceSceneData.Triangles.size() > 0)
1516   {
1517     anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceVertexBuffer, CL_FALSE,
1518                                      0, myRaytraceSceneData.Vertices.size() * sizeof(cl_float4),
1519                                      &myRaytraceSceneData.Vertices.front(),
1520                                      0, NULL, NULL);
1521     anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNormalBuffer, CL_FALSE,
1522                                      0, myRaytraceSceneData.Normals.size() * sizeof(cl_float4),
1523                                      &myRaytraceSceneData.Normals.front(),
1524                                      0, NULL, NULL);
1525     anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceTriangleBuffer, CL_FALSE,
1526                                      0, myRaytraceSceneData.Triangles.size() * sizeof(cl_int4),
1527                                      &myRaytraceSceneData.Triangles.front(),
1528                                      0, NULL, NULL);
1529     if (anError != CL_SUCCESS)
1530     {
1531   #ifdef RAY_TRACE_PRINT_INFO
1532       std::cout << "Error! Failed to write OpenCL geometry buffers!" << std::endl;
1533   #endif
1534       return Standard_False;
1535     }
1536   }
1537
1538   // Write BVH buffers
1539   if (aTree.DataRcrdBuffer().size() > 0)
1540   {
1541     anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeMinPointBuffer, CL_FALSE,
1542                                      0, aTree.MinPointBuffer().size() * sizeof(cl_float4),
1543                                      &aTree.MinPointBuffer().front(),
1544                                      0, NULL, NULL);
1545     anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeMaxPointBuffer, CL_FALSE,
1546                                      0, aTree.MaxPointBuffer().size() * sizeof(cl_float4),
1547                                      &aTree.MaxPointBuffer().front(),
1548                                      0, NULL, NULL);
1549     anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeDataRcrdBuffer, CL_FALSE,
1550                                      0, aTree.DataRcrdBuffer().size() * sizeof(cl_int4),
1551                                      &aTree.DataRcrdBuffer().front(),
1552                                      0, NULL, NULL);
1553     if (anError != CL_SUCCESS)
1554     {
1555   #ifdef RAY_TRACE_PRINT_INFO
1556       std::cout << "Error! Failed to write OpenCL BVH buffers!" << std::endl;
1557   #endif
1558       return Standard_False;
1559     }
1560   }
1561
1562   // Write material buffers
1563   if (myRaytraceSceneData.Materials.size() > 0)
1564   {
1565     const size_t aSize    = myRaytraceSceneData.Materials.size();
1566     const void*  aDataPtr = myRaytraceSceneData.Materials.front().Packed();
1567
1568     anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceMaterialBuffer, CL_FALSE,
1569                                      0, aSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr,
1570                                      0, NULL, NULL);
1571     if (anError != CL_SUCCESS)
1572     {
1573   #ifdef RAY_TRACE_PRINT_INFO
1574       std::cout << "Error! Failed to write OpenCL material buffer!" << std::endl;
1575   #endif
1576       return Standard_False;
1577     }
1578   }
1579
1580   anError |= clFinish (myRaytraceQueue);
1581 #ifdef RAY_TRACE_PRINT_INFO
1582   if (anError != CL_SUCCESS)
1583     std::cout << "Error! Failed to set scene data buffers!" << std::endl;
1584 #endif
1585
1586   if (anError == CL_SUCCESS)
1587     myIsRaytraceDataValid = myRaytraceSceneData.Triangles.size() > 0;
1588
1589 #ifdef RAY_TRACE_PRINT_INFO
1590
1591   float aMemUsed = static_cast<float> (
1592     myRaytraceSceneData.Materials.size()) * sizeof (OpenGl_RaytraceMaterial);
1593
1594   aMemUsed += static_cast<float> (
1595     myRaytraceSceneData.Triangles.size() * sizeof (OpenGl_RTVec4i) +
1596     myRaytraceSceneData.Vertices.size()  * sizeof (OpenGl_RTVec4f) +
1597     myRaytraceSceneData.Normals.size()   * sizeof (OpenGl_RTVec4f));
1598
1599   aMemUsed += static_cast<float> (
1600     aTree.MinPointBuffer().size() * sizeof (OpenGl_RTVec4f) +
1601     aTree.MaxPointBuffer().size() * sizeof (OpenGl_RTVec4f) +
1602     aTree.DataRcrdBuffer().size() * sizeof (OpenGl_RTVec4i));
1603
1604   std::cout << "GPU memory used (Mb): " << aMemUsed / 1e6f << std::endl;
1605
1606 #endif
1607
1608   myRaytraceSceneData.Clear();
1609
1610   myBVHBuilder.CleanUp();
1611
1612   return (CL_SUCCESS == anError);
1613 }
1614
1615 #define OPENCL_GROUP_SIZE_TEST_
1616
1617 // =======================================================================
1618 // function : RunRaytraceOpenCLKernels
1619 // purpose  : Runs OpenCL ray-tracing kernels
1620 // =======================================================================
1621 Standard_Boolean OpenGl_Workspace::RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView,
1622                                                              const GLfloat          theOrigins[16],
1623                                                              const GLfloat          theDirects[16],
1624                                                              const Standard_Integer theSizeX,
1625                                                              const Standard_Integer theSizeY)
1626 {
1627   if (myRaytraceRenderKernel == NULL || myRaytraceQueue == NULL)
1628     return Standard_False;
1629
1630   ////////////////////////////////////////////////////////////
1631   // Set kernel arguments
1632
1633   cl_uint anIndex = 0;
1634   cl_int  anError = 0;
1635
1636   anError  = clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1637                              sizeof(cl_mem), &myRaytraceOutputImage);
1638   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1639                              sizeof(cl_mem), &myRaytraceEnvironment);
1640   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1641                              sizeof(cl_mem), &myRaytraceNodeMinPointBuffer);
1642   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1643                              sizeof(cl_mem), &myRaytraceNodeMaxPointBuffer);
1644   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1645                              sizeof(cl_mem), &myRaytraceNodeDataRcrdBuffer);
1646   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1647                              sizeof(cl_mem), &myRaytraceLightSourceBuffer);
1648   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1649                              sizeof(cl_mem), &myRaytraceMaterialBuffer);
1650   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1651                              sizeof(cl_mem), &myRaytraceVertexBuffer);
1652   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1653                              sizeof(cl_mem), &myRaytraceNormalBuffer);
1654   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1655                              sizeof(cl_mem), &myRaytraceTriangleBuffer);
1656
1657   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1658                              sizeof(cl_float16), theOrigins);
1659   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1660                              sizeof(cl_float16), theDirects);
1661
1662   cl_int aLightCount =  static_cast<cl_int> (myRaytraceSceneData.LightSources.size());
1663
1664   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1665                              sizeof(cl_int),   &aLightCount);
1666   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1667                              sizeof(cl_float), &myRaytraceSceneEpsilon);
1668   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1669                              sizeof(cl_float), &myRaytraceSceneRadius);
1670   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1671                              sizeof(cl_int),   &theCView.IsShadowsEnabled);
1672   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1673                              sizeof(cl_int),   &theCView.IsReflectionsEnabled);
1674   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1675                              sizeof(cl_int),   &theSizeX);
1676   anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
1677                              sizeof(cl_int),   &theSizeY);
1678   if (anError != CL_SUCCESS)
1679   {
1680     const TCollection_ExtendedString aMsg = "Error! Failed to set arguments of ray-tracing kernel!";
1681     myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1682                               GL_DEBUG_TYPE_ERROR_ARB,
1683                               0,
1684                               GL_DEBUG_SEVERITY_HIGH_ARB,
1685                               aMsg);
1686     return Standard_False;
1687   }
1688
1689   // Note: second-pass 'smoothing' kernel runs only if anti-aliasing is enabled
1690   if (theCView.IsAntialiasingEnabled)
1691   {
1692     anIndex = 0;
1693     anError  = clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1694                                sizeof(cl_mem), &myRaytraceOutputImage);
1695     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1696                                sizeof(cl_mem), &myRaytraceOutputImageAA);
1697     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1698                                sizeof(cl_mem), &myRaytraceEnvironment);
1699     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1700                                sizeof(cl_mem), &myRaytraceNodeMinPointBuffer);
1701     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1702                                sizeof(cl_mem), &myRaytraceNodeMaxPointBuffer);
1703     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1704                                sizeof(cl_mem), &myRaytraceNodeDataRcrdBuffer);
1705     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1706                                sizeof(cl_mem), &myRaytraceLightSourceBuffer);
1707     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1708                                sizeof(cl_mem), &myRaytraceMaterialBuffer);
1709     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1710                                sizeof(cl_mem), &myRaytraceVertexBuffer);
1711     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1712                                sizeof(cl_mem), &myRaytraceNormalBuffer);
1713     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1714                                sizeof(cl_mem), &myRaytraceTriangleBuffer);
1715
1716     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1717                                sizeof(cl_float16), theOrigins);
1718     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1719                                 sizeof(cl_float16), theDirects);
1720
1721     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1722                                sizeof(cl_int),   &aLightCount);
1723     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1724                                sizeof(cl_float), &myRaytraceSceneEpsilon);
1725     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1726                                sizeof(cl_float), &myRaytraceSceneRadius);
1727     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1728                                sizeof(cl_int),   &theCView.IsShadowsEnabled);
1729     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1730                                sizeof(cl_int),   &theCView.IsReflectionsEnabled);
1731     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1732                                sizeof(cl_int),   &theSizeX);
1733     anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
1734                                sizeof(cl_int),   &theSizeY);
1735     if (anError != CL_SUCCESS)
1736     {
1737       const TCollection_ExtendedString aMsg = "Error! Failed to set arguments of 'smoothing' kernel!";
1738       myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1739                                 GL_DEBUG_TYPE_ERROR_ARB,
1740                                 0,
1741                                 GL_DEBUG_SEVERITY_HIGH_ARB,
1742                                 aMsg);
1743       return Standard_False;
1744     }
1745   }
1746
1747   // Set work size
1748   size_t aLocSizeRender[] = { myIsAmdComputePlatform ? 2 : 4, 32 };
1749
1750 #ifdef OPENCL_GROUP_SIZE_TEST
1751   for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1 )
1752   for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1 )
1753 #endif
1754   {
1755 #ifdef OPENCL_GROUP_SIZE_TEST
1756     aLocSizeRender[0] = aLocX;
1757     aLocSizeRender[1] = aLocY;
1758 #endif
1759
1760     size_t aWorkSizeX = theSizeX;
1761     if (aWorkSizeX % aLocSizeRender[0] != 0)
1762       aWorkSizeX += aLocSizeRender[0] - aWorkSizeX % aLocSizeRender[0];
1763
1764     size_t aWokrSizeY = theSizeY;
1765     if (aWokrSizeY % aLocSizeRender[1] != 0 )
1766       aWokrSizeY += aLocSizeRender[1] - aWokrSizeY % aLocSizeRender[1];
1767
1768     size_t aGlbSizeRender[] = { aWorkSizeX, aWokrSizeY };
1769
1770     // Run kernel
1771     cl_event anEvent (NULL), anEventSmooth (NULL);
1772     anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceRenderKernel,
1773                                       2, NULL, aGlbSizeRender, aLocSizeRender,
1774                                       0, NULL, &anEvent);
1775     if (anError != CL_SUCCESS)
1776     {
1777       const TCollection_ExtendedString aMsg = "Error! Failed to execute the ray-tracing kernel!";
1778       myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1779                                 GL_DEBUG_TYPE_ERROR_ARB,
1780                                 0,
1781                                 GL_DEBUG_SEVERITY_HIGH_ARB,
1782                                 aMsg);
1783       return Standard_False;
1784     }
1785     clWaitForEvents (1, &anEvent);
1786
1787     if (theCView.IsAntialiasingEnabled)
1788     {
1789       size_t aLocSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4,
1790                                   myIsAmdComputePlatform ? 8 : 32 };
1791
1792 #ifdef OPENCL_GROUP_SIZE_TEST
1793       aLocSizeSmooth[0] = aLocX;
1794       aLocSizeSmooth[1] = aLocY;
1795 #endif
1796
1797       aWorkSizeX = theSizeX;
1798       if (aWorkSizeX % aLocSizeSmooth[0] != 0)
1799         aWorkSizeX += aLocSizeSmooth[0] - aWorkSizeX % aLocSizeSmooth[0];
1800
1801       size_t aWokrSizeY = theSizeY;
1802       if (aWokrSizeY % aLocSizeSmooth[1] != 0 )
1803         aWokrSizeY += aLocSizeSmooth[1] - aWokrSizeY % aLocSizeSmooth[1];
1804
1805       size_t aGlbSizeSmooth [] = { aWorkSizeX, aWokrSizeY };
1806       anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceSmoothKernel,
1807                                         2, NULL, aGlbSizeSmooth, aLocSizeSmooth,
1808                                         0, NULL, &anEventSmooth);
1809       clWaitForEvents (1, &anEventSmooth);
1810
1811       if (anError != CL_SUCCESS)
1812       {
1813         const TCollection_ExtendedString aMsg = "Error! Failed to execute the 'smoothing' kernel!";
1814         myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1815                                   GL_DEBUG_TYPE_ERROR_ARB,
1816                                   0,
1817                                   GL_DEBUG_SEVERITY_HIGH_ARB,
1818                                   aMsg);
1819         return Standard_False;
1820       }
1821     }
1822
1823     // Get the profiling data
1824 #if defined (RAY_TRACE_PRINT_INFO) || defined(OPENCL_GROUP_SIZE_TEST)
1825
1826     cl_ulong aTimeStart,
1827              aTimeFinal;
1828
1829     clGetEventProfilingInfo (anEvent, CL_PROFILING_COMMAND_START,
1830                              sizeof(aTimeStart), &aTimeStart, NULL);
1831     clGetEventProfilingInfo (anEvent, CL_PROFILING_COMMAND_END,
1832                              sizeof(aTimeFinal), &aTimeFinal, NULL);
1833     std::cout << "\tRender time (ms): " << ( aTimeFinal - aTimeStart ) / 1e6f << std::endl;
1834
1835     if (theCView.IsAntialiasingEnabled)
1836     {
1837       clGetEventProfilingInfo (anEventSmooth, CL_PROFILING_COMMAND_START,
1838                                sizeof(aTimeStart), &aTimeStart, NULL);
1839       clGetEventProfilingInfo (anEventSmooth, CL_PROFILING_COMMAND_END,
1840                                sizeof(aTimeFinal), &aTimeFinal, NULL);
1841       std::cout << "\tSmoothing time (ms): " << ( aTimeFinal - aTimeStart ) / 1e6f << std::endl;
1842     }
1843 #endif
1844
1845     if (anEvent != NULL)
1846       clReleaseEvent (anEvent);
1847
1848     if (anEventSmooth != NULL)
1849       clReleaseEvent (anEventSmooth);
1850   }
1851
1852   return Standard_True;
1853 }
1854
1855 // =======================================================================
1856 // function : ComputeInverseMatrix
1857 // purpose  : Computes inversion of 4x4 floating-point matrix
1858 // =======================================================================
1859 template <typename T>
1860 void ComputeInverseMatrix (const T m[16], T inv[16])
1861 {
1862   inv[ 0] = m[ 5] * (m[10] * m[15] - m[11] * m[14]) -
1863             m[ 9] * (m[ 6] * m[15] - m[ 7] * m[14]) -
1864             m[13] * (m[ 7] * m[10] - m[ 6] * m[11]);
1865
1866   inv[ 1] = m[ 1] * (m[11] * m[14] - m[10] * m[15]) -
1867             m[ 9] * (m[ 3] * m[14] - m[ 2] * m[15]) -
1868             m[13] * (m[ 2] * m[11] - m[ 3] * m[10]);
1869
1870   inv[ 2] = m[ 1] * (m[ 6] * m[15] - m[ 7] * m[14]) -
1871             m[ 5] * (m[ 2] * m[15] - m[ 3] * m[14]) -
1872             m[13] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]);
1873
1874   inv[ 3] = m[ 1] * (m[ 7] * m[10] - m[ 6] * m[11]) -
1875             m[ 5] * (m[ 3] * m[10] - m[ 2] * m[11]) -
1876             m[ 9] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]);
1877
1878   inv[ 4] = m[ 4] * (m[11] * m[14] - m[10] * m[15]) -
1879             m[ 8] * (m[ 7] * m[14] - m[ 6] * m[15]) -
1880             m[12] * (m[ 6] * m[11] - m[ 7] * m[10]);
1881
1882   inv[ 5] = m[ 0] * (m[10] * m[15] - m[11] * m[14]) -
1883             m[ 8] * (m[ 2] * m[15] - m[ 3] * m[14]) -
1884             m[12] * (m[ 3] * m[10] - m[ 2] * m[11]);
1885
1886   inv[ 6] = m[ 0] * (m[ 7] * m[14] - m[ 6] * m[15]) -
1887             m[ 4] * (m[ 3] * m[14] - m[ 2] * m[15]) -
1888             m[12] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]);
1889
1890   inv[ 7] = m[ 0] * (m[ 6] * m[11] - m[ 7] * m[10]) -
1891             m[ 4] * (m[ 2] * m[11] - m[ 3] * m[10]) -
1892             m[ 8] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]);
1893
1894   inv[ 8] = m[ 4] * (m[ 9] * m[15] - m[11] * m[13]) -
1895             m[ 8] * (m[ 5] * m[15] - m[ 7] * m[13]) -
1896             m[12] * (m[ 7] * m[ 9] - m[ 5] * m[11]);
1897
1898   inv[ 9] = m[ 0] * (m[11] * m[13] - m[ 9] * m[15]) -
1899             m[ 8] * (m[ 3] * m[13] - m[ 1] * m[15]) -
1900             m[12] * (m[ 1] * m[11] - m[ 3] * m[ 9]);
1901
1902   inv[10] = m[ 0] * (m[ 5] * m[15] - m[ 7] * m[13]) -
1903             m[ 4] * (m[ 1] * m[15] - m[ 3] * m[13]) -
1904             m[12] * (m[ 3] * m[ 5] - m[ 1] * m[ 7]);
1905
1906   inv[11] = m[ 0] * (m[ 7] * m[ 9] - m[ 5] * m[11]) -
1907             m[ 4] * (m[ 3] * m[ 9] - m[ 1] * m[11]) -
1908             m[ 8] * (m[ 1] * m[ 7] - m[ 3] * m[ 5]);
1909
1910   inv[12] = m[ 4] * (m[10] * m[13] - m[ 9] * m[14]) -
1911             m[ 8] * (m[ 6] * m[13] - m[ 5] * m[14]) -
1912             m[12] * (m[ 5] * m[10] - m[ 6] * m[ 9]);
1913
1914   inv[13] = m[ 0] * (m[ 9] * m[14] - m[10] * m[13]) -
1915             m[ 8] * (m[ 1] * m[14] - m[ 2] * m[13]) -
1916             m[12] * (m[ 2] * m[ 9] - m[ 1] * m[10]);
1917
1918   inv[14] = m[ 0] * (m[ 6] * m[13] - m[ 5] * m[14]) -
1919             m[ 4] * (m[ 2] * m[13] - m[ 1] * m[14]) -
1920             m[12] * (m[ 1] * m[ 6] - m[ 2] * m[ 5]);
1921
1922   inv[15] = m[ 0] * (m[ 5] * m[10] - m[ 6] * m[ 9]) -
1923             m[ 4] * (m[ 1] * m[10] - m[ 2] * m[ 9]) -
1924             m[ 8] * (m[ 2] * m[ 5] - m[ 1] * m[ 6]);
1925
1926   T det = m[0] * inv[ 0] +
1927           m[1] * inv[ 4] +
1928           m[2] * inv[ 8] +
1929           m[3] * inv[12];
1930
1931   if (det == T (0.0)) return;
1932
1933   det = T (1.0) / det;
1934
1935   for (Standard_Integer i = 0; i < 16; ++i)
1936     inv[i] *= det;
1937 }
1938
1939 // =======================================================================
1940 // function : GenerateCornerRays
1941 // purpose  : Generates primary rays for corners of screen quad
1942 // =======================================================================
1943 void GenerateCornerRays (const GLdouble theInvModelProj[16],
1944                          cl_float       theOrigins[16],
1945                          cl_float       theDirects[16])
1946 {
1947   Standard_Integer aOriginIndex = 0;
1948   Standard_Integer aDirectIndex = 0;
1949
1950   for (Standard_Integer y = -1; y <= 1; y += 2)
1951   {
1952     for (Standard_Integer x = -1; x <= 1; x += 2)
1953     {
1954       OpenGl_RTVec4f aOrigin (float(x),
1955                               float(y),
1956                               -1.f,
1957                               1.f);
1958
1959       aOrigin = MatVecMult (theInvModelProj, aOrigin);
1960
1961       OpenGl_RTVec4f aDirect (float(x),
1962                               float(y),
1963                               1.f,
1964                               1.f);
1965
1966       aDirect = MatVecMult (theInvModelProj, aDirect) - aOrigin;
1967
1968       GLdouble aInvLen = 1.f / sqrt (aDirect.x() * aDirect.x() +
1969                                      aDirect.y() * aDirect.y() +
1970                                      aDirect.z() * aDirect.z());
1971
1972       theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.x());
1973       theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.y());
1974       theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.z());
1975       theOrigins [aOriginIndex++] = 1.f;
1976
1977       theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.x() * aInvLen);
1978       theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.y() * aInvLen);
1979       theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.z() * aInvLen);
1980       theDirects [aDirectIndex++] = 0.f;
1981     }
1982   }
1983 }
1984
1985 // =======================================================================
1986 // function : Raytrace
1987 // purpose  : Redraws the window using OpenCL ray tracing
1988 // =======================================================================
1989 Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView,
1990                                              const Standard_Integer theSizeX,
1991                                              const Standard_Integer theSizeY,
1992                                              const Tint             theToSwap)
1993 {
1994   if (!InitOpenCL())
1995     return Standard_False;
1996
1997   if (!ResizeRaytraceOutputBuffer (theSizeX, theSizeY))
1998     return Standard_False;
1999
2000   if (!UpdateRaytraceEnvironmentMap())
2001     return Standard_False;
2002
2003   if (!UpdateRaytraceGeometry (Standard_True))
2004     return Standard_False;
2005
2006   // Get model-view and projection matrices
2007   TColStd_Array2OfReal theOrientation (0, 3, 0, 3);
2008   TColStd_Array2OfReal theViewMapping (0, 3, 0, 3);
2009
2010   myView->GetMatrices (theOrientation, theViewMapping, Standard_True);
2011
2012   GLdouble aOrientationMatrix[16];
2013   GLdouble aViewMappingMatrix[16];
2014   GLdouble aOrientationInvers[16];
2015
2016   for (Standard_Integer j = 0; j < 4; ++j)
2017     for (Standard_Integer i = 0; i < 4; ++i)
2018     {
2019       aOrientationMatrix [4 * j + i] = theOrientation (i, j);
2020       aViewMappingMatrix [4 * j + i] = theViewMapping (i, j);
2021     }
2022
2023   ComputeInverseMatrix (aOrientationMatrix, aOrientationInvers);
2024
2025   if (!UpdateRaytraceLightSources (aOrientationInvers))
2026     return Standard_False;
2027
2028   // Generate primary rays for corners of the screen quad
2029   glMatrixMode (GL_MODELVIEW);
2030
2031   glLoadMatrixd (aViewMappingMatrix);
2032   glMultMatrixd (aOrientationMatrix);
2033
2034   GLdouble aModelProject[16];
2035   GLdouble aInvModelProj[16];
2036
2037   glGetDoublev (GL_MODELVIEW_MATRIX, aModelProject);
2038
2039   ComputeInverseMatrix (aModelProject, aInvModelProj);
2040
2041   GLfloat aOrigins[16];
2042   GLfloat aDirects[16];
2043
2044   GenerateCornerRays (aInvModelProj,
2045                       aOrigins,
2046                       aDirects);
2047
2048   // Compute ray-traced image using OpenCL kernel
2049   cl_mem anImages[] = { myRaytraceOutputImage, myRaytraceOutputImageAA };
2050   cl_int anError = clEnqueueAcquireGLObjects (myRaytraceQueue,
2051                                               2, anImages,
2052                                               0, NULL, NULL);
2053   clFinish (myRaytraceQueue);
2054
2055   if (myIsRaytraceDataValid)
2056   {
2057     RunRaytraceOpenCLKernels (theCView,
2058                               aOrigins,
2059                               aDirects,
2060                               theSizeX,
2061                               theSizeY);
2062   }
2063
2064   anError |= clEnqueueReleaseGLObjects (myRaytraceQueue,
2065                                         2, anImages,
2066                                         0, NULL, NULL);
2067   clFinish (myRaytraceQueue);
2068
2069   // Draw background
2070   glPushAttrib (GL_ENABLE_BIT |
2071                 GL_CURRENT_BIT |
2072                 GL_COLOR_BUFFER_BIT |
2073                 GL_DEPTH_BUFFER_BIT);
2074
2075   glDisable (GL_DEPTH_TEST);
2076
2077   if (NamedStatus & OPENGL_NS_WHITEBACK)
2078   {
2079     glClearColor (1.0f, 1.0f, 1.0f, 1.0f);
2080   }
2081   else
2082   {
2083     glClearColor (myBgColor.rgb[0],
2084                   myBgColor.rgb[1],
2085                   myBgColor.rgb[2],
2086                   1.0f);
2087   }
2088
2089   glClear (GL_COLOR_BUFFER_BIT);
2090
2091   Handle(OpenGl_Workspace) aWorkspace (this);
2092   myView->DrawBackground (aWorkspace);
2093
2094   // Draw dummy quad to show result image
2095   glEnable (GL_COLOR_MATERIAL);
2096   glEnable (GL_BLEND);
2097
2098   glDisable (GL_DEPTH_TEST);
2099
2100   glBlendFunc (GL_ONE, GL_SRC_ALPHA);
2101
2102   glEnable (GL_TEXTURE_RECTANGLE);
2103
2104   glMatrixMode (GL_PROJECTION);
2105   glLoadIdentity();
2106
2107   glMatrixMode (GL_MODELVIEW);
2108   glLoadIdentity();
2109
2110   glColor3f (1.0f, 1.0f, 1.0f);
2111
2112   if (!theCView.IsAntialiasingEnabled)
2113     myRaytraceOutputTexture->Bind (myGlContext);
2114   else
2115     myRaytraceOutputTextureAA->Bind (myGlContext);
2116
2117   if (myIsRaytraceDataValid)
2118   {
2119     glBegin (GL_QUADS);
2120     {
2121       glTexCoord2i (       0,        0);   glVertex2f (-1.f, -1.f);
2122       glTexCoord2i (       0, theSizeY);   glVertex2f (-1.f,  1.f);
2123       glTexCoord2i (theSizeX, theSizeY);   glVertex2f ( 1.f,  1.f);
2124       glTexCoord2i (theSizeX,        0);   glVertex2f ( 1.f, -1.f);
2125     }
2126     glEnd();
2127   }
2128
2129   glPopAttrib();
2130
2131   // Swap the buffers
2132   if (theToSwap)
2133   {
2134     GetGlContext()->SwapBuffers();
2135     myBackBufferRestored = Standard_False;
2136   }
2137   else
2138     glFlush();
2139
2140   return Standard_True;
2141 }
2142
2143 #endif