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