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