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