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