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