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