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