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