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