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