1 // Created on: 2013-08-27
2 // Created by: Denis BOGOLEPOV
3 // Copyright (c) 2013 OPEN CASCADE SAS
5 // This file is part of Open CASCADE Technology software library.
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
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.
13 // Alternatively, this file may be used under the terms of Open CASCADE
14 // commercial license or contractual agreement.
22 #include <OpenGl_Cl.hxx>
29 #pragma comment (lib, "DelayImp.lib")
30 #pragma comment (lib, "OpenCL.lib")
32 #elif defined(__APPLE__) && !defined(MACOSX_USE_GLX)
33 #include <OpenGL/CGLCurrent.h>
38 #include <OpenGl_Context.hxx>
39 #include <OpenGl_Texture.hxx>
40 #include <OpenGl_View.hxx>
41 #include <OpenGl_Workspace.hxx>
42 #include <Standard_Assert.hxx>
44 using namespace OpenGl_Raytrace;
46 //! Use this macro to output ray-tracing debug info
47 //#define RAY_TRACE_PRINT_INFO
49 #ifdef RAY_TRACE_PRINT_INFO
50 #include <OSD_Timer.hxx>
53 //! OpenCL source of ray-tracing kernels.
54 extern const char THE_RAY_TRACE_OPENCL_SOURCE[];
56 // =======================================================================
57 // function : MatVecMult
58 // purpose : Multiples 4x4 matrix by 4D vector
59 // =======================================================================
61 BVH_Vec4f MatVecMult (const T m[16], const BVH_Vec4f& v)
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()));
74 // =======================================================================
75 // function : UpdateRaytraceEnvironmentMap
76 // purpose : Updates environment map for ray-tracing
77 // =======================================================================
78 Standard_Boolean OpenGl_Workspace::UpdateRaytraceEnvironmentMap()
81 return Standard_False;
83 if (myViewModificationStatus == myView->ModificationState())
86 cl_int anError = CL_SUCCESS;
88 if (myRaytraceEnvironment != NULL)
89 clReleaseMemObject (myRaytraceEnvironment);
91 Standard_Integer aSizeX = 1;
92 Standard_Integer aSizeY = 1;
94 if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE)
96 aSizeX = (myView->TextureEnv()->SizeX() <= 0) ? 1 : myView->TextureEnv()->SizeX();
97 aSizeY = (myView->TextureEnv()->SizeY() <= 0) ? 1 : myView->TextureEnv()->SizeY();
100 cl_image_format anImageFormat;
102 anImageFormat.image_channel_order = CL_RGBA;
103 anImageFormat.image_channel_data_type = CL_FLOAT;
105 myRaytraceEnvironment = clCreateImage2D (myComputeContext,
106 CL_MEM_READ_ONLY, &anImageFormat, aSizeX, aSizeY, 0, NULL, &anError);
108 cl_float* aPixelData = new cl_float[aSizeX * aSizeY * 4];
110 // Note: texture format is not compatible with OpenCL image
111 // (it's not possible to create image directly from texture)
113 if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE)
115 myView->TextureEnv()->Bind (GetGlContext());
117 glGetTexImage (GL_TEXTURE_2D,
123 myView->TextureEnv()->Unbind (GetGlContext());
127 for (Standard_Integer aPixel = 0; aPixel < aSizeX * aSizeY * 4; ++aPixel)
128 aPixelData[aPixel] = 0.f;
131 size_t anImageOffset[] = { 0,
135 size_t anImageRegion[] = { aSizeX,
139 anError |= clEnqueueWriteImage (myComputeQueue, myRaytraceEnvironment,
140 CL_TRUE, anImageOffset, anImageRegion, 0, 0, aPixelData, 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;
149 myViewModificationStatus = myView->ModificationState();
151 return (anError == CL_SUCCESS);
154 // =======================================================================
155 // function : UpdateRaytraceGeometry
156 // purpose : Updates 3D scene geometry for ray tracing
157 // =======================================================================
158 Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theCheck)
161 return Standard_False;
163 // Note: In 'check' mode the scene geometry is analyzed for modifications
164 // This is light-weight procedure performed for each frame
168 myRaytraceGeometry.Clear();
170 myIsRaytraceDataValid = Standard_False;
174 if (myLayersModificationStatus != myView->LayerList().ModificationState())
176 return UpdateRaytraceGeometry (Standard_False);
180 Standard_ShortReal* aTransform (NULL);
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;
187 const OpenGl_LayerList& aList = myView->LayerList();
189 for (OpenGl_SequenceOfLayers::Iterator anLayerIt (aList.Layers()); anLayerIt.More(); anLayerIt.Next())
191 const OpenGl_PriorityList& aPriorityList = anLayerIt.Value();
193 if (aPriorityList.NbStructures() == 0)
196 const OpenGl_ArrayOfStructure& aStructArray = aPriorityList.ArrayOfStructures();
198 for (Standard_Integer anIndex = 0; anIndex < aStructArray.Length(); ++anIndex)
200 OpenGl_SequenceOfStructure::Iterator aStructIt;
202 for (aStructIt.Init (aStructArray (anIndex)); aStructIt.More(); aStructIt.Next())
204 const OpenGl_Structure* aStructure = aStructIt.Value();
208 if (CheckRaytraceStructure (aStructure))
210 return UpdateRaytraceGeometry (Standard_False);
215 if (!aStructure->IsRaytracable())
218 if (aStructure->Transformation()->mat != NULL)
220 if (aTransform == NULL)
221 aTransform = new Standard_ShortReal[16];
223 for (Standard_Integer i = 0; i < 4; ++i)
224 for (Standard_Integer j = 0; j < 4; ++j)
226 aTransform[j * 4 + i] = aStructure->Transformation()->mat[i][j];
230 AddRaytraceStructure (aStructure, aTransform, anElements);
238 // Actualize the hash map of structures -- remove out-of-date records
239 std::map<const OpenGl_Structure*, Standard_Size>::iterator anIter = myStructureStates.begin();
241 while (anIter != myStructureStates.end())
243 if (anElements.find (anIter->first) == anElements.end())
245 myStructureStates.erase (anIter++);
253 // Actualize OpenGL layer list state
254 myLayersModificationStatus = myView->LayerList().ModificationState();
256 // Rebuild bottom-level and high-level BVHs
257 myRaytraceGeometry.ProcessAcceleration();
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())));
264 myRaytraceSceneRadius = 2.f /* scale factor */ * Max (aMinRadius, aMaxRadius);
266 myRaytraceSceneEpsilon = Max (1e-4f,
267 myRaytraceGeometry.Box().Size().Length() * 1e-4f);
269 return WriteRaytraceSceneToDevice();
272 delete [] aTransform;
274 return Standard_True;
277 // =======================================================================
278 // function : CheckRaytraceStructure
279 // purpose : Checks to see if the structure is modified
280 // =======================================================================
281 Standard_Boolean OpenGl_Workspace::CheckRaytraceStructure (const OpenGl_Structure* theStructure)
283 if (!theStructure->IsRaytracable())
285 // Checks to see if all ray-tracable elements were
286 // removed from the structure
287 if (theStructure->ModificationState() > 0)
289 theStructure->ResetModificationState();
290 return Standard_True;
293 return Standard_False;
296 std::map<const OpenGl_Structure*, Standard_Size>::iterator aStructState = myStructureStates.find (theStructure);
298 if (aStructState != myStructureStates.end())
299 return aStructState->second != theStructure->ModificationState();
301 return Standard_True;
304 // =======================================================================
305 // function : CreateMaterial
306 // purpose : Creates ray-tracing material properties
307 // =======================================================================
308 void CreateMaterial (const OPENGL_SURF_PROP& theProp, OpenGl_RaytraceMaterial& theMaterial)
310 const float* aSrcAmb = theProp.isphysic ? theProp.ambcol.rgb : theProp.matcol.rgb;
311 theMaterial.Ambient = BVH_Vec4f (aSrcAmb[0] * theProp.amb,
312 aSrcAmb[1] * theProp.amb,
313 aSrcAmb[2] * theProp.amb,
316 const float* aSrcDif = theProp.isphysic ? theProp.difcol.rgb : theProp.matcol.rgb;
317 theMaterial.Diffuse = BVH_Vec4f (aSrcDif[0] * theProp.diff,
318 aSrcDif[1] * theProp.diff,
319 aSrcDif[2] * theProp.diff,
322 const float aDefSpecCol[4] = {1.0f, 1.0f, 1.0f, 1.0f};
323 const float* aSrcSpe = theProp.isphysic ? theProp.speccol.rgb : aDefSpecCol;
324 theMaterial.Specular = BVH_Vec4f (aSrcSpe[0] * theProp.spec,
325 aSrcSpe[1] * theProp.spec,
326 aSrcSpe[2] * theProp.spec,
329 const float* aSrcEms = theProp.isphysic ? theProp.emscol.rgb : theProp.matcol.rgb;
330 theMaterial.Emission = BVH_Vec4f (aSrcEms[0] * theProp.emsv,
331 aSrcEms[1] * theProp.emsv,
332 aSrcEms[2] * theProp.emsv,
335 // Note: Here we use sub-linear transparency function
336 // to produce realistic-looking transparency effect
337 theMaterial.Transparency = BVH_Vec4f (powf (theProp.trans, 0.75f),
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()));
346 const float aReflectionScale = 0.75f / aMaxRefl;
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;
354 // =======================================================================
355 // function : AddRaytraceStructure
356 // purpose : Adds OpenGL structure to ray-traced scene geometry
357 // =======================================================================
358 Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure* theStructure,
359 const Standard_ShortReal* theTransform, std::set<const OpenGl_Structure*>& theElements)
361 theElements.insert (theStructure);
363 if (!theStructure->IsVisible())
365 myStructureStates[theStructure] = theStructure->ModificationState();
366 return Standard_True;
369 // Get structure material
370 Standard_Integer aStructMatID = -1;
372 if (theStructure->AspectFace() != NULL)
374 aStructMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
376 OpenGl_RaytraceMaterial aStructMaterial;
377 CreateMaterial (theStructure->AspectFace()->IntFront(), aStructMaterial);
379 myRaytraceGeometry.Materials.push_back (aStructMaterial);
382 for (OpenGl_ListOfGroup::Iterator anItg (theStructure->Groups()); anItg.More(); anItg.Next())
384 // Get group material
385 Standard_Integer aGroupMatID = -1;
387 if (anItg.Value()->AspectFace() != NULL)
389 aGroupMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
391 OpenGl_RaytraceMaterial aGroupMaterial;
392 CreateMaterial (anItg.Value()->AspectFace()->IntFront(), aGroupMaterial);
394 myRaytraceGeometry.Materials.push_back (aGroupMaterial);
397 Standard_Integer aMatID = aGroupMatID < 0 ? aStructMatID : aGroupMatID;
401 aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
403 myRaytraceGeometry.Materials.push_back (OpenGl_RaytraceMaterial());
406 // Add OpenGL elements from group (extract primitives arrays and aspects)
407 for (const OpenGl_ElementNode* aNode = anItg.Value()->FirstNode(); aNode != NULL; aNode = aNode->next)
409 OpenGl_AspectFace* anAspect = dynamic_cast<OpenGl_AspectFace*> (aNode->elem);
410 if (anAspect != NULL)
412 aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
414 OpenGl_RaytraceMaterial aMaterial;
415 CreateMaterial (anAspect->IntFront(), aMaterial);
417 myRaytraceGeometry.Materials.push_back (aMaterial);
421 OpenGl_PrimitiveArray* aPrimArray = dynamic_cast<OpenGl_PrimitiveArray*> (aNode->elem);
422 if (aPrimArray != NULL)
424 NCollection_Handle<BVH_Object<Standard_ShortReal, 4> > aSet =
425 AddRaytracePrimitiveArray (aPrimArray->PArray(), aMatID, theTransform);
428 myRaytraceGeometry.Objects().Append (aSet);
434 Standard_ShortReal* aTransform (NULL);
436 // Process all connected OpenGL structures
437 for (OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures()); anIts.More(); anIts.Next())
439 if (anIts.Value()->Transformation()->mat != NULL)
441 Standard_ShortReal* aTransform = new Standard_ShortReal[16];
443 for (Standard_Integer i = 0; i < 4; ++i)
444 for (Standard_Integer j = 0; j < 4; ++j)
446 aTransform[j * 4 + i] =
447 anIts.Value()->Transformation()->mat[i][j];
451 if (anIts.Value()->IsRaytracable())
452 AddRaytraceStructure (anIts.Value(), aTransform != NULL ? aTransform : theTransform, theElements);
457 myStructureStates[theStructure] = theStructure->ModificationState();
459 return Standard_True;
462 // =======================================================================
463 // function : AddRaytracePrimitiveArray
464 // purpose : Adds OpenGL primitive array to ray-traced scene geometry
465 // =======================================================================
466 OpenGl_TriangleSet* OpenGl_Workspace::AddRaytracePrimitiveArray (
467 const CALL_DEF_PARRAY* theArray, Standard_Integer theMatID, const Standard_ShortReal* theTransform)
469 if (theArray->type != TelPolygonsArrayType &&
470 theArray->type != TelTrianglesArrayType &&
471 theArray->type != TelQuadranglesArrayType &&
472 theArray->type != TelTriangleFansArrayType &&
473 theArray->type != TelTriangleStripsArrayType &&
474 theArray->type != TelQuadrangleStripsArrayType)
479 if (theArray->vertices == NULL)
482 #ifdef RAY_TRACE_PRINT_INFO
483 switch (theArray->type)
485 case TelPolygonsArrayType:
486 std::cout << "\tAdding TelPolygonsArrayType" << std::endl; break;
487 case TelTrianglesArrayType:
488 std::cout << "\tAdding TelTrianglesArrayType" << std::endl; break;
489 case TelQuadranglesArrayType:
490 std::cout << "\tAdding TelQuadranglesArrayType" << std::endl; break;
491 case TelTriangleFansArrayType:
492 std::cout << "\tAdding TelTriangleFansArrayType" << std::endl; break;
493 case TelTriangleStripsArrayType:
494 std::cout << "\tAdding TelTriangleStripsArrayType" << std::endl; break;
495 case TelQuadrangleStripsArrayType:
496 std::cout << "\tAdding TelQuadrangleStripsArrayType" << std::endl; break;
500 OpenGl_TriangleSet* aSet = new OpenGl_TriangleSet;
503 aSet->Vertices.reserve (theArray->num_vertexs);
505 for (Standard_Integer aVert = 0; aVert < theArray->num_vertexs; ++aVert)
507 BVH_Vec4f aVertex (theArray->vertices[aVert].xyz[0],
508 theArray->vertices[aVert].xyz[1],
509 theArray->vertices[aVert].xyz[2],
512 aVertex = MatVecMult (theTransform, aVertex);
514 aSet->Vertices.push_back (aVertex);
517 aSet->Normals.reserve (theArray->num_vertexs);
519 for (Standard_Integer aNorm = 0; aNorm < theArray->num_vertexs; ++aNorm)
523 // Note: In case of absence of normals, the
524 // renderer uses generated geometric normals
526 if (theArray->vnormals != NULL)
528 aNormal = BVH_Vec4f (theArray->vnormals[aNorm].xyz[0],
529 theArray->vnormals[aNorm].xyz[1],
530 theArray->vnormals[aNorm].xyz[2],
534 aNormal = MatVecMult (theTransform, aNormal);
537 aSet->Normals.push_back (aNormal);
540 if (theArray->num_bounds > 0)
542 #ifdef RAY_TRACE_PRINT_INFO
543 std::cout << "\tNumber of bounds = " << theArray->num_bounds << std::endl;
546 Standard_Integer aBoundStart = 0;
548 for (Standard_Integer aBound = 0; aBound < theArray->num_bounds; ++aBound)
550 const Standard_Integer aVertNum = theArray->bounds[aBound];
552 #ifdef RAY_TRACE_PRINT_INFO
553 std::cout << "\tAdding indices from bound " << aBound << ": " <<
554 aBoundStart << " .. " << aVertNum << std::endl;
557 if (!AddRaytraceVertexIndices (aSet, theArray, aBoundStart, aVertNum, theMatID))
563 aBoundStart += aVertNum;
568 const Standard_Integer aVertNum = theArray->num_edges > 0 ? theArray->num_edges : theArray->num_vertexs;
570 #ifdef RAY_TRACE_PRINT_INFO
571 std::cout << "\tAdding indices from array: " << aVertNum << std::endl;
574 if (!AddRaytraceVertexIndices (aSet, theArray, 0, aVertNum, theMatID))
582 if (aSet->Size() != 0)
588 // =======================================================================
589 // function : AddRaytraceVertexIndices
590 // purpose : Adds vertex indices to ray-traced scene geometry
591 // =======================================================================
592 Standard_Boolean OpenGl_Workspace::AddRaytraceVertexIndices (OpenGl_TriangleSet* theSet,
593 const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
595 switch (theArray->type)
597 case TelTrianglesArrayType:
598 return AddRaytraceTriangleArray (theSet, theArray, theOffset, theCount, theMatID);
600 case TelQuadranglesArrayType:
601 return AddRaytraceQuadrangleArray (theSet, theArray, theOffset, theCount, theMatID);
603 case TelTriangleFansArrayType:
604 return AddRaytraceTriangleFanArray (theSet, theArray, theOffset, theCount, theMatID);
606 case TelTriangleStripsArrayType:
607 return AddRaytraceTriangleStripArray (theSet, theArray, theOffset, theCount, theMatID);
609 case TelQuadrangleStripsArrayType:
610 return AddRaytraceQuadrangleStripArray (theSet, theArray, theOffset, theCount, theMatID);
613 return AddRaytracePolygonArray (theSet, theArray, theOffset, theCount, theMatID);
617 // =======================================================================
618 // function : AddRaytraceTriangleArray
619 // purpose : Adds OpenGL triangle array to ray-traced scene geometry
620 // =======================================================================
621 Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleArray (OpenGl_TriangleSet* theSet,
622 const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
625 return Standard_True;
627 theSet->Elements.reserve (theSet->Elements.size() + theCount / 3);
629 if (theArray->num_edges > 0)
631 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; aVert += 3)
633 theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
634 theArray->edges[aVert + 1],
635 theArray->edges[aVert + 2],
641 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; aVert += 3)
643 theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
650 return Standard_True;
653 // =======================================================================
654 // function : AddRaytraceTriangleFanArray
655 // purpose : Adds OpenGL triangle fan array to ray-traced scene geometry
656 // =======================================================================
657 Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleFanArray (OpenGl_TriangleSet* theSet,
658 const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
661 return Standard_True;
663 theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
665 if (theArray->num_edges > 0)
667 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
669 theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset],
670 theArray->edges[aVert + 1],
671 theArray->edges[aVert + 2],
677 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
679 theSet->Elements.push_back (BVH_Vec4i (theOffset,
686 return Standard_True;
689 // =======================================================================
690 // function : AddRaytraceTriangleStripArray
691 // purpose : Adds OpenGL triangle strip array to ray-traced scene geometry
692 // =======================================================================
693 Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleStripArray (OpenGl_TriangleSet* theSet,
694 const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
697 return Standard_True;
699 theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
701 if (theArray->num_edges > 0)
703 for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2)
705 theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + aCW ? 1 : 0],
706 theArray->edges[aVert + aCW ? 0 : 1],
707 theArray->edges[aVert + 2],
713 for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2)
715 theSet->Elements.push_back (BVH_Vec4i (aVert + aCW ? 1 : 0,
722 return Standard_True;
725 // =======================================================================
726 // function : AddRaytraceQuadrangleArray
727 // purpose : Adds OpenGL quad array to ray-traced scene geometry
728 // =======================================================================
729 Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleArray (OpenGl_TriangleSet* theSet,
730 const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
733 return Standard_True;
735 theSet->Elements.reserve (theSet->Elements.size() + theCount / 2);
737 if (theArray->num_edges > 0)
739 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 4)
741 theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
742 theArray->edges[aVert + 1],
743 theArray->edges[aVert + 2],
746 theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
747 theArray->edges[aVert + 2],
748 theArray->edges[aVert + 3],
754 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 4)
756 theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
761 theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
768 return Standard_True;
771 // =======================================================================
772 // function : AddRaytraceQuadrangleStripArray
773 // purpose : Adds OpenGL quad strip array to ray-traced scene geometry
774 // =======================================================================
775 Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleStripArray (OpenGl_TriangleSet* theSet,
776 const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
779 return Standard_True;
781 theSet->Elements.reserve (theSet->Elements.size() + 2 * theCount - 6);
783 if (theArray->num_edges > 0)
785 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 2)
787 theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
788 theArray->edges[aVert + 1],
789 theArray->edges[aVert + 2],
792 theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 1],
793 theArray->edges[aVert + 3],
794 theArray->edges[aVert + 2],
800 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 2)
802 theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
807 theSet->Elements.push_back (BVH_Vec4i (aVert + 1,
814 return Standard_True;
817 // =======================================================================
818 // function : AddRaytracePolygonArray
819 // purpose : Adds OpenGL polygon array to ray-traced scene geometry
820 // =======================================================================
821 Standard_Boolean OpenGl_Workspace::AddRaytracePolygonArray (OpenGl_TriangleSet* theSet,
822 const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
825 return Standard_True;
827 theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
829 if (theArray->num_edges > 0)
831 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
833 theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset],
834 theArray->edges[aVert + 1],
835 theArray->edges[aVert + 2],
841 for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
843 theSet->Elements.push_back (BVH_Vec4i (theOffset,
850 return Standard_True;
853 // =======================================================================
854 // function : UpdateRaytraceLightSources
855 // purpose : Updates 3D scene light sources for ray-tracing
856 // =======================================================================
857 Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble theInvModelView[16])
859 myRaytraceGeometry.Sources.clear();
861 myRaytraceGeometry.GlobalAmbient = BVH_Vec4f (0.0f, 0.0f, 0.0f, 0.0f);
863 for (OpenGl_ListOfLight::Iterator anItl (myView->LightList()); anItl.More(); anItl.Next())
865 const OpenGl_Light& aLight = anItl.Value();
867 if (aLight.Type == Visual3d_TOLS_AMBIENT)
869 myRaytraceGeometry.GlobalAmbient += BVH_Vec4f (aLight.Color.r(),
876 BVH_Vec4f aDiffuse (aLight.Color.r(),
881 BVH_Vec4f aPosition (-aLight.Direction.x(),
882 -aLight.Direction.y(),
883 -aLight.Direction.z(),
886 if (aLight.Type != Visual3d_TOLS_DIRECTIONAL)
888 aPosition = BVH_Vec4f (aLight.Position.x(),
894 if (aLight.IsHeadlight)
895 aPosition = MatVecMult (theInvModelView, aPosition);
897 myRaytraceGeometry.Sources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition));
900 cl_int anError = CL_SUCCESS;
902 if (myRaytraceLightSourceBuffer != NULL)
903 clReleaseMemObject (myRaytraceLightSourceBuffer);
905 Standard_Integer aLightBufferSize = myRaytraceGeometry.Sources.size() != 0 ?
906 static_cast<Standard_Integer> (myRaytraceGeometry.Sources.size()) : 1;
908 myRaytraceLightSourceBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
909 aLightBufferSize * sizeof(OpenGl_RaytraceLight), NULL, &anError);
911 if (myRaytraceGeometry.Sources.size() != 0)
913 const void* aDataPtr = myRaytraceGeometry.Sources.front().Packed();
915 anError |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0,
916 aLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr, 0, NULL, NULL);
919 #ifdef RAY_TRACE_PRINT_INFO
920 if (anError != CL_SUCCESS)
922 std::cout << "Error! Failed to set light sources";
924 return Standard_False;
928 return Standard_True;
931 // =======================================================================
932 // function : CheckOpenCL
933 // purpose : Checks OpenCL dynamic library availability
934 // =======================================================================
935 Standard_Boolean CheckOpenCL()
937 #if defined ( _WIN32 )
941 cl_uint aNbPlatforms;
942 clGetPlatformIDs (0, NULL, &aNbPlatforms);
944 __except (EXCEPTION_EXECUTE_HANDLER)
946 return Standard_False;
951 return Standard_True;
954 // =======================================================================
955 // function : InitOpenCL
956 // purpose : Initializes OpenCL objects
957 // =======================================================================
958 Standard_Boolean OpenGl_Workspace::InitOpenCL()
960 if (myComputeInitStatus != OpenGl_CLIS_NONE)
962 return myComputeInitStatus == OpenGl_CLIS_INIT;
967 myComputeInitStatus = OpenGl_CLIS_FAIL; // fail to load OpenCL library
968 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
969 GL_DEBUG_TYPE_ERROR_ARB,
971 GL_DEBUG_SEVERITY_HIGH_ARB,
972 "Failed to load OpenCL dynamic library!");
973 return Standard_False;
976 // Obtain the list of platforms available
977 cl_uint aNbPlatforms = 0;
978 cl_int anError = clGetPlatformIDs (0, NULL, &aNbPlatforms);
979 cl_platform_id* aPlatforms = (cl_platform_id* )alloca (aNbPlatforms * sizeof(cl_platform_id));
980 anError |= clGetPlatformIDs (aNbPlatforms, aPlatforms, NULL);
981 if (anError != CL_SUCCESS
982 || aNbPlatforms == 0)
984 myComputeInitStatus = OpenGl_CLIS_FAIL;
985 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
986 GL_DEBUG_TYPE_ERROR_ARB,
988 GL_DEBUG_SEVERITY_HIGH_ARB,
989 "No any OpenCL platform installed!");
990 return Standard_False;
993 // Note: We try to find NVIDIA or AMD platforms with GPU devices!
994 cl_platform_id aPrefPlatform = NULL;
995 for (cl_uint aPlatIter = 0; aPlatIter < aNbPlatforms; ++aPlatIter)
998 anError = clGetPlatformInfo (aPlatforms[aPlatIter], CL_PLATFORM_NAME,
999 sizeof(aName), aName, NULL);
1000 if (anError != CL_SUCCESS)
1005 if (strncmp (aName, "NVIDIA", strlen ("NVIDIA")) == 0)
1007 aPrefPlatform = aPlatforms[aPlatIter];
1009 // Use optimizations for NVIDIA GPUs
1010 myIsAmdComputePlatform = Standard_False;
1012 else if (strncmp (aName, "AMD", strlen ("AMD")) == 0)
1014 aPrefPlatform = (aPrefPlatform == NULL)
1015 ? aPlatforms[aPlatIter]
1018 // Use optimizations for ATI/AMD platform
1019 myIsAmdComputePlatform = Standard_True;
1023 if (aPrefPlatform == NULL)
1025 aPrefPlatform = aPlatforms[0];
1028 // Obtain the list of devices available in the selected platform
1029 cl_uint aNbDevices = 0;
1030 anError = clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU,
1031 0, NULL, &aNbDevices);
1033 cl_device_id* aDevices = (cl_device_id* )alloca (aNbDevices * sizeof(cl_device_id));
1034 anError |= clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU,
1035 aNbDevices, aDevices, NULL);
1036 if (anError != CL_SUCCESS)
1038 myComputeInitStatus = OpenGl_CLIS_FAIL;
1039 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1040 GL_DEBUG_TYPE_ERROR_ARB,
1042 GL_DEBUG_SEVERITY_HIGH_ARB,
1043 "Failed to get OpenCL GPU device!");
1044 return Standard_False;
1047 // Note: Simply get first available GPU
1048 cl_device_id aDevice = aDevices[0];
1050 // detect old contexts
1051 char aVerClStr[256];
1052 clGetDeviceInfo (aDevice, CL_DEVICE_VERSION,
1053 sizeof(aVerClStr), aVerClStr, NULL);
1054 aVerClStr[strlen ("OpenCL 1.0")] = '\0';
1055 const bool isVer10 = strncmp (aVerClStr, "OpenCL 1.0", strlen ("OpenCL 1.0")) == 0;
1057 // Create OpenCL context
1058 cl_context_properties aCtxProp[] =
1060 #if defined(__APPLE__) && !defined(MACOSX_USE_GLX)
1061 CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE,
1062 (cl_context_properties )CGLGetShareGroup (CGLGetCurrentContext()),
1063 #elif defined(_WIN32)
1064 CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform,
1065 CL_GL_CONTEXT_KHR, (cl_context_properties )wglGetCurrentContext(),
1066 CL_WGL_HDC_KHR, (cl_context_properties )wglGetCurrentDC(),
1068 CL_GL_CONTEXT_KHR, (cl_context_properties )glXGetCurrentContext(),
1069 CL_GLX_DISPLAY_KHR, (cl_context_properties )glXGetCurrentDisplay(),
1070 CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform,
1075 myComputeContext = clCreateContext (aCtxProp,
1076 #if defined(__APPLE__) && !defined(MACOSX_USE_GLX)
1077 0, NULL, // device will be taken from GL context
1081 NULL, NULL, &anError);
1082 if (anError != CL_SUCCESS)
1084 myComputeInitStatus = OpenGl_CLIS_FAIL;
1085 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1086 GL_DEBUG_TYPE_ERROR_ARB,
1088 GL_DEBUG_SEVERITY_HIGH_ARB,
1089 "Failed to initialize OpenCL context!");
1090 return Standard_False;
1093 // Create OpenCL program
1094 const char* aSources[] =
1096 isVer10 ? "#define M_PI_F ( float )( 3.14159265359f )\n" : "",
1097 THE_RAY_TRACE_OPENCL_SOURCE
1099 myRaytraceProgram = clCreateProgramWithSource (myComputeContext, 2,
1100 aSources, NULL, &anError);
1101 if (anError != CL_SUCCESS)
1103 myComputeInitStatus = OpenGl_CLIS_FAIL;
1104 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1105 GL_DEBUG_TYPE_ERROR_ARB,
1107 GL_DEBUG_SEVERITY_HIGH_ARB,
1108 "Failed to create OpenCL ray-tracing program!");
1109 return Standard_False;
1112 anError = clBuildProgram (myRaytraceProgram, 0,
1113 NULL, NULL, NULL, NULL);
1117 cl_int aResult = clGetProgramBuildInfo (myRaytraceProgram, aDevice,
1118 CL_PROGRAM_BUILD_LOG, 0, NULL, &aLogLen);
1120 char* aBuildLog = (char* )alloca (aLogLen);
1121 aResult |= clGetProgramBuildInfo (myRaytraceProgram, aDevice,
1122 CL_PROGRAM_BUILD_LOG, aLogLen, aBuildLog, NULL);
1123 if (aResult == CL_SUCCESS)
1125 if (anError != CL_SUCCESS)
1127 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1128 GL_DEBUG_TYPE_ERROR_ARB,
1130 GL_DEBUG_SEVERITY_HIGH_ARB,
1135 #ifdef RAY_TRACE_PRINT_INFO
1136 std::cout << aBuildLog << std::endl;
1142 if (anError != CL_SUCCESS)
1144 return Standard_False;
1147 // Create OpenCL ray tracing kernels
1148 myRaytraceRenderKernel = clCreateKernel (myRaytraceProgram, "RaytraceRender", &anError);
1149 if (anError != CL_SUCCESS)
1151 myComputeInitStatus = OpenGl_CLIS_FAIL;
1152 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1153 GL_DEBUG_TYPE_ERROR_ARB,
1155 GL_DEBUG_SEVERITY_HIGH_ARB,
1156 "Failed to create OpenCL ray-tracing kernel!");
1157 return Standard_False;
1160 myRaytraceSmoothKernel = clCreateKernel (myRaytraceProgram, "RaytraceSmooth", &anError);
1161 if (anError != CL_SUCCESS)
1163 myComputeInitStatus = OpenGl_CLIS_FAIL;
1164 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1165 GL_DEBUG_TYPE_ERROR_ARB,
1167 GL_DEBUG_SEVERITY_HIGH_ARB,
1168 "Failed to create OpenCL ray-tracing kernel!");
1169 return Standard_False;
1172 // Create OpenCL command queue
1173 // Note: For profiling set CL_QUEUE_PROFILING_ENABLE
1174 cl_command_queue_properties aProps = CL_QUEUE_PROFILING_ENABLE;
1176 myComputeQueue = clCreateCommandQueue (myComputeContext, aDevice, aProps, &anError);
1177 if (anError != CL_SUCCESS)
1179 myComputeInitStatus = OpenGl_CLIS_FAIL;
1180 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1181 GL_DEBUG_TYPE_ERROR_ARB,
1183 GL_DEBUG_SEVERITY_HIGH_ARB,
1184 "Failed to create OpenCL command queue!");
1186 return Standard_False;
1189 myComputeInitStatus = OpenGl_CLIS_INIT; // initialized in normal way
1190 return Standard_True;
1193 // =======================================================================
1194 // function : GetOpenClDeviceInfo
1195 // purpose : Returns information about device used for computations
1196 // =======================================================================
1197 Standard_Boolean OpenGl_Workspace::GetOpenClDeviceInfo (NCollection_DataMap<TCollection_AsciiString,
1198 TCollection_AsciiString>& theInfo) const
1201 if (myComputeContext == NULL)
1203 return Standard_False;
1206 size_t aDevicesSize = 0;
1207 cl_int anError = clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, 0, NULL, &aDevicesSize);
1208 cl_device_id* aDevices = (cl_device_id* )alloca (aDevicesSize);
1209 anError |= clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, aDevicesSize, aDevices, NULL);
1210 if (anError != CL_SUCCESS)
1212 return Standard_False;
1215 char aDeviceName[256];
1216 anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_NAME, sizeof(aDeviceName), aDeviceName, NULL);
1217 theInfo.Bind ("Name", aDeviceName);
1219 char aDeviceVendor[256];
1220 anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_VENDOR, sizeof(aDeviceVendor), aDeviceVendor, NULL);
1221 theInfo.Bind ("Vendor", aDeviceVendor);
1223 cl_device_type aDeviceType;
1224 anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_TYPE, sizeof(aDeviceType), &aDeviceType, NULL);
1225 theInfo.Bind ("Type", aDeviceType == CL_DEVICE_TYPE_GPU ? "GPU" : "CPU");
1226 return Standard_True;
1229 // =======================================================================
1230 // function : ReleaseOpenCL
1231 // purpose : Releases resources of OpenCL objects
1232 // =======================================================================
1233 void OpenGl_Workspace::ReleaseOpenCL()
1235 clReleaseKernel (myRaytraceRenderKernel);
1236 clReleaseKernel (myRaytraceSmoothKernel);
1238 clReleaseProgram (myRaytraceProgram);
1239 clReleaseCommandQueue (myComputeQueue);
1241 clReleaseMemObject (myRaytraceOutputImage);
1242 clReleaseMemObject (myRaytraceEnvironment);
1243 clReleaseMemObject (myRaytraceOutputImageAA);
1245 clReleaseMemObject (myRaytraceMaterialBuffer);
1246 clReleaseMemObject (myRaytraceLightSourceBuffer);
1248 clReleaseMemObject (mySceneNodeInfoBuffer);
1249 clReleaseMemObject (mySceneMinPointBuffer);
1250 clReleaseMemObject (mySceneMaxPointBuffer);
1252 clReleaseMemObject (myObjectNodeInfoBuffer);
1253 clReleaseMemObject (myObjectMinPointBuffer);
1254 clReleaseMemObject (myObjectMaxPointBuffer);
1256 clReleaseMemObject (myGeometryVertexBuffer);
1257 clReleaseMemObject (myGeometryNormalBuffer);
1258 clReleaseMemObject (myGeometryTriangBuffer);
1260 clReleaseContext (myComputeContext);
1262 if (!myGlContext.IsNull())
1264 if (!myRaytraceOutputTexture.IsNull())
1265 myGlContext->DelayedRelease (myRaytraceOutputTexture);
1266 myRaytraceOutputTexture.Nullify();
1268 if (!myRaytraceOutputTextureAA.IsNull())
1269 myGlContext->DelayedRelease (myRaytraceOutputTextureAA);
1270 myRaytraceOutputTextureAA.Nullify();
1274 // =======================================================================
1275 // function : ResizeRaytraceOutputBuffer
1276 // purpose : Resizes OpenCL output image
1277 // =======================================================================
1278 Standard_Boolean OpenGl_Workspace::ResizeRaytraceOutputBuffer (const cl_int theSizeX,
1279 const cl_int theSizeY)
1281 if (myComputeContext == NULL)
1283 return Standard_False;
1286 if (!myRaytraceOutputTexture.IsNull())
1288 Standard_Boolean toResize = myRaytraceOutputTexture->SizeX() != theSizeX ||
1289 myRaytraceOutputTexture->SizeY() != theSizeY;
1292 return Standard_True;
1294 if (!myGlContext.IsNull())
1296 if (!myRaytraceOutputTexture.IsNull())
1297 myGlContext->DelayedRelease (myRaytraceOutputTexture);
1298 if (!myRaytraceOutputTextureAA.IsNull())
1299 myGlContext->DelayedRelease (myRaytraceOutputTextureAA);
1303 myRaytraceOutputTexture = new OpenGl_Texture();
1305 myRaytraceOutputTexture->Create (myGlContext);
1306 myRaytraceOutputTexture->InitRectangle (myGlContext,
1307 theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>());
1309 myRaytraceOutputTextureAA = new OpenGl_Texture();
1311 myRaytraceOutputTextureAA->Create (myGlContext);
1312 myRaytraceOutputTextureAA->InitRectangle (myGlContext,
1313 theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>());
1315 if (myRaytraceOutputImage != NULL)
1316 clReleaseMemObject (myRaytraceOutputImage);
1318 if (myRaytraceOutputImageAA != NULL)
1319 clReleaseMemObject (myRaytraceOutputImageAA);
1321 cl_int anError = CL_SUCCESS;
1323 myRaytraceOutputImage = clCreateFromGLTexture2D (myComputeContext,
1324 CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTexture->TextureId(), &anError);
1326 if (anError != CL_SUCCESS)
1328 #ifdef RAY_TRACE_PRINT_INFO
1329 std::cout << "Error! Failed to create output image!" << std::endl;
1331 return Standard_False;
1334 myRaytraceOutputImageAA = clCreateFromGLTexture2D (myComputeContext,
1335 CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTextureAA->TextureId(), &anError);
1337 if (anError != CL_SUCCESS)
1339 #ifdef RAY_TRACE_PRINT_INFO
1340 std::cout << "Error! Failed to create anti-aliased output image!" << std::endl;
1342 return Standard_False;
1345 return Standard_True;
1348 // =======================================================================
1349 // function : WriteRaytraceSceneToDevice
1350 // purpose : Writes scene geometry to OpenCL device
1351 // =======================================================================
1352 Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice()
1354 if (myComputeContext == NULL)
1355 return Standard_False;
1357 cl_int anErrorRes = CL_SUCCESS;
1359 if (mySceneNodeInfoBuffer != NULL)
1360 anErrorRes |= clReleaseMemObject (mySceneNodeInfoBuffer);
1362 if (mySceneMinPointBuffer != NULL)
1363 anErrorRes |= clReleaseMemObject (mySceneMinPointBuffer);
1365 if (mySceneMaxPointBuffer != NULL)
1366 anErrorRes |= clReleaseMemObject (mySceneMaxPointBuffer);
1368 if (myObjectNodeInfoBuffer != NULL)
1369 anErrorRes |= clReleaseMemObject (myObjectNodeInfoBuffer);
1371 if (myObjectMinPointBuffer != NULL)
1372 anErrorRes |= clReleaseMemObject (myObjectMinPointBuffer);
1374 if (myObjectMaxPointBuffer != NULL)
1375 anErrorRes |= clReleaseMemObject (myObjectMaxPointBuffer);
1377 if (myGeometryVertexBuffer != NULL)
1378 anErrorRes |= clReleaseMemObject (myGeometryVertexBuffer);
1380 if (myGeometryNormalBuffer != NULL)
1381 anErrorRes |= clReleaseMemObject (myGeometryNormalBuffer);
1383 if (myGeometryTriangBuffer != NULL)
1384 anErrorRes |= clReleaseMemObject (myGeometryTriangBuffer);
1386 if (myRaytraceMaterialBuffer != NULL)
1387 anErrorRes |= clReleaseMemObject (myRaytraceMaterialBuffer);
1389 if (anErrorRes != CL_SUCCESS)
1391 #ifdef RAY_TRACE_PRINT_INFO
1392 std::cout << "Error! Failed to release OpenCL buffers" << std::endl;
1394 return Standard_False;
1397 /////////////////////////////////////////////////////////////////////////////
1398 // Create material buffer
1400 const size_t aMaterialBufferSize =
1401 myRaytraceGeometry.Materials.size() != 0 ? myRaytraceGeometry.Materials.size() : 1;
1403 myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext,
1404 CL_MEM_READ_ONLY, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL, &anErrorRes);
1406 if (anErrorRes != CL_SUCCESS)
1408 #ifdef RAY_TRACE_PRINT_INFO
1409 std::cout << "Error! Failed to create OpenCL material buffer" << std::endl;
1411 return Standard_False;
1414 /////////////////////////////////////////////////////////////////////////////
1415 // Create BVHs buffers
1417 cl_int anErrorTmp = CL_SUCCESS;
1419 const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = myRaytraceGeometry.BVH();
1421 const size_t aSceneMinPointBufferSize =
1422 aBVH->MinPointBuffer().size() != 0 ? aBVH->MinPointBuffer().size() : 1;
1424 mySceneMinPointBuffer = clCreateBuffer (myComputeContext,
1425 CL_MEM_READ_ONLY, aSceneMinPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp);
1426 anErrorRes |= anErrorTmp;
1428 const size_t aSceneMaxPointBufferSize =
1429 aBVH->MaxPointBuffer().size() != 0 ? aBVH->MaxPointBuffer().size() : 1;
1431 mySceneMaxPointBuffer = clCreateBuffer (myComputeContext,
1432 CL_MEM_READ_ONLY, aSceneMaxPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp);
1433 anErrorRes |= anErrorTmp;
1435 const size_t aSceneNodeInfoBufferSize =
1436 aBVH->NodeInfoBuffer().size() != 0 ? aBVH->NodeInfoBuffer().size() : 1;
1438 mySceneNodeInfoBuffer = clCreateBuffer (myComputeContext,
1439 CL_MEM_READ_ONLY, aSceneNodeInfoBufferSize * sizeof(cl_int4), NULL, &anErrorTmp);
1440 anErrorRes |= anErrorTmp;
1442 if (anErrorRes != CL_SUCCESS)
1444 #ifdef RAY_TRACE_PRINT_INFO
1445 std::cout << "Error! Failed to create OpenCL buffers for high-level scene BVH" << std::endl;
1447 return Standard_False;
1450 Standard_Integer aTotalVerticesNb = 0;
1451 Standard_Integer aTotalElementsNb = 0;
1452 Standard_Integer aTotalBVHNodesNb = 0;
1454 for (Standard_Integer anElemIndex = 0; anElemIndex < myRaytraceGeometry.Size(); ++anElemIndex)
1456 OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
1457 myRaytraceGeometry.Objects().ChangeValue (anElemIndex).operator->());
1459 Standard_ASSERT_RETURN (aTriangleSet != NULL,
1460 "Error! Failed to get triangulation of OpenGL element", Standard_False);
1462 aTotalVerticesNb += (int)aTriangleSet->Vertices.size();
1463 aTotalElementsNb += (int)aTriangleSet->Elements.size();
1465 Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(),
1466 "Error! Failed to get bottom-level BVH of OpenGL element", Standard_False);
1468 aTotalBVHNodesNb += (int)aTriangleSet->BVH()->NodeInfoBuffer().size();
1471 aTotalBVHNodesNb = aTotalBVHNodesNb > 0 ? aTotalBVHNodesNb : 1;
1473 myObjectNodeInfoBuffer = clCreateBuffer (myComputeContext,
1474 CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_int4), NULL, &anErrorTmp);
1475 anErrorRes |= anErrorTmp;
1477 myObjectMinPointBuffer = clCreateBuffer (myComputeContext,
1478 CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp);
1479 anErrorRes |= anErrorTmp;
1481 myObjectMaxPointBuffer = clCreateBuffer (myComputeContext,
1482 CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp);
1483 anErrorRes |= anErrorTmp;
1485 if (anErrorRes != CL_SUCCESS)
1487 #ifdef RAY_TRACE_PRINT_INFO
1488 std::cout << "Error! Failed to create OpenCL buffers for bottom-level scene BVHs" << std::endl;
1490 return Standard_False;
1493 /////////////////////////////////////////////////////////////////////////////
1494 // Create geometry buffers
1496 aTotalVerticesNb = aTotalVerticesNb > 0 ? aTotalVerticesNb : 1;
1498 myGeometryVertexBuffer = clCreateBuffer (myComputeContext,
1499 CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp);
1500 anErrorRes |= anErrorTmp;
1502 myGeometryNormalBuffer = clCreateBuffer (myComputeContext,
1503 CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp);
1504 anErrorRes |= anErrorTmp;
1506 aTotalElementsNb = aTotalElementsNb > 0 ? aTotalElementsNb : 1;
1508 myGeometryTriangBuffer = clCreateBuffer (myComputeContext,
1509 CL_MEM_READ_ONLY, aTotalElementsNb * sizeof(cl_int4), NULL, &anErrorTmp);
1510 anErrorRes |= anErrorTmp;
1512 if (anErrorRes != CL_SUCCESS)
1514 #ifdef RAY_TRACE_PRINT_INFO
1515 std::cout << "Error! Failed to create OpenCL geometry buffers" << std::endl;
1517 return Standard_False;
1520 /////////////////////////////////////////////////////////////////////////////
1521 // Write BVH and geometry buffers
1523 if (aBVH->NodeInfoBuffer().size() != 0)
1525 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneNodeInfoBuffer, CL_FALSE, 0,
1526 aSceneNodeInfoBufferSize * sizeof(cl_int4), &aBVH->NodeInfoBuffer().front(), 0, NULL, NULL);
1528 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMinPointBuffer, CL_FALSE, 0,
1529 aSceneMinPointBufferSize * sizeof(cl_float4), &aBVH->MinPointBuffer().front(), 0, NULL, NULL);
1531 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMaxPointBuffer, CL_FALSE, 0,
1532 aSceneMaxPointBufferSize * sizeof(cl_float4), &aBVH->MaxPointBuffer().front(), 0, NULL, NULL);
1534 anErrorRes |= clFinish (myComputeQueue);
1536 if (anErrorRes != CL_SUCCESS)
1538 #ifdef RAY_TRACE_PRINT_INFO
1539 std::cout << "Error! Failed to write OpenCL buffers for high-level scene BVH" << std::endl;
1541 return Standard_False;
1544 for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx)
1546 if (!aBVH->IsOuter (aNodeIdx))
1549 OpenGl_TriangleSet* aTriangleSet = myRaytraceGeometry.TriangleSet (aNodeIdx);
1551 Standard_ASSERT_RETURN (aTriangleSet != NULL,
1552 "Error! Failed to get triangulation of OpenGL element", Standard_False);
1554 const size_t aBVHBuffserSize =
1555 aTriangleSet->BVH()->NodeInfoBuffer().size() != 0 ? aTriangleSet->BVH()->NodeInfoBuffer().size() : 1;
1557 const Standard_Integer aBVHOffset = myRaytraceGeometry.AccelerationOffset (aNodeIdx);
1559 Standard_ASSERT_RETURN (aBVHOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
1560 "Error! Failed to get offset for bottom-level BVH", Standard_False);
1562 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectNodeInfoBuffer, CL_FALSE, aBVHOffset * sizeof(cl_int4),
1563 aBVHBuffserSize * sizeof(cl_int4), &aTriangleSet->BVH()->NodeInfoBuffer().front(), 0, NULL, NULL);
1565 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMinPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4),
1566 aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MinPointBuffer().front(), 0, NULL, NULL);
1568 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMaxPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4),
1569 aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MaxPointBuffer().front(), 0, NULL, NULL);
1571 anErrorRes |= clFinish (myComputeQueue);
1573 if (anErrorRes != CL_SUCCESS)
1575 #ifdef RAY_TRACE_PRINT_INFO
1576 std::cout << "Error! Failed to write OpenCL buffers for bottom-level scene BVHs" << std::endl;
1578 return Standard_False;
1581 const Standard_Integer aVerticesOffset = myRaytraceGeometry.VerticesOffset (aNodeIdx);
1583 Standard_ASSERT_RETURN (aVerticesOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
1584 "Error! Failed to get offset for triangulation vertices of OpenGL element", Standard_False);
1586 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryVertexBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4),
1587 aTriangleSet->Vertices.size() * sizeof(cl_float4), &aTriangleSet->Vertices.front(), 0, NULL, NULL);
1589 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryNormalBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4),
1590 aTriangleSet->Normals.size() * sizeof(cl_float4), &aTriangleSet->Normals.front(), 0, NULL, NULL);
1592 const Standard_Integer anElementsOffset = myRaytraceGeometry.ElementsOffset (aNodeIdx);
1594 Standard_ASSERT_RETURN (anElementsOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
1595 "Error! Failed to get offset for triangulation elements of OpenGL element", Standard_False);
1597 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryTriangBuffer, CL_FALSE, anElementsOffset * sizeof(cl_int4),
1598 aTriangleSet->Elements.size() * sizeof(cl_int4), &aTriangleSet->Elements.front(), 0, NULL, NULL);
1600 anErrorRes |= clFinish (myComputeQueue);
1602 if (anErrorRes != CL_SUCCESS)
1604 #ifdef RAY_TRACE_PRINT_INFO
1605 std::cout << "Error! Failed to write OpenCL triangulation buffers for OpenGL element" << std::endl;
1607 return Standard_False;
1612 /////////////////////////////////////////////////////////////////////////////
1613 // Write material buffer
1615 if (myRaytraceGeometry.Materials.size() != 0)
1617 const void* aDataPtr = myRaytraceGeometry.Materials.front().Packed();
1619 anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceMaterialBuffer,
1620 CL_FALSE, 0, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr, 0, NULL, NULL);
1622 if (anErrorRes != CL_SUCCESS)
1624 #ifdef RAY_TRACE_PRINT_INFO
1625 std::cout << "Error! Failed to write OpenCL material buffer" << std::endl;
1627 return Standard_False;
1631 anErrorRes |= clFinish (myComputeQueue);
1633 if (anErrorRes == CL_SUCCESS)
1635 myIsRaytraceDataValid = myRaytraceGeometry.Objects().Size() != 0;
1637 #ifdef RAY_TRACE_PRINT_INFO
1640 std::cout << "Error! Failed to set scene data buffers" << std::endl;
1644 #ifdef RAY_TRACE_PRINT_INFO
1646 Standard_ShortReal aMemUsed = 0.f;
1648 for (Standard_Integer anElemIdx = 0; anElemIdx < myRaytraceGeometry.Size(); ++anElemIdx)
1650 OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
1651 myRaytraceGeometry.Objects().ChangeValue (anElemIdx).operator->());
1653 aMemUsed += static_cast<Standard_ShortReal> (
1654 aTriangleSet->Vertices.size() * sizeof (BVH_Vec4f));
1655 aMemUsed += static_cast<Standard_ShortReal> (
1656 aTriangleSet->Normals.size() * sizeof (BVH_Vec4f));
1657 aMemUsed += static_cast<Standard_ShortReal> (
1658 aTriangleSet->Elements.size() * sizeof (BVH_Vec4i));
1660 aMemUsed += static_cast<Standard_ShortReal> (
1661 aTriangleSet->BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
1662 aMemUsed += static_cast<Standard_ShortReal> (
1663 aTriangleSet->BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
1664 aMemUsed += static_cast<Standard_ShortReal> (
1665 aTriangleSet->BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
1668 aMemUsed += static_cast<Standard_ShortReal> (
1669 myRaytraceGeometry.BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
1670 aMemUsed += static_cast<Standard_ShortReal> (
1671 myRaytraceGeometry.BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
1672 aMemUsed += static_cast<Standard_ShortReal> (
1673 myRaytraceGeometry.BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
1675 std::cout << "GPU Memory Used (MB): ~" << aMemUsed / 1048576 << std::endl;
1679 return (CL_SUCCESS == anErrorRes);
1682 // Use it to estimate the optimal size of OpenCL work group
1683 // #define OPENCL_GROUP_SIZE_TEST
1685 // =======================================================================
1686 // function : RunRaytraceOpenCLKernels
1687 // purpose : Runs OpenCL ray-tracing kernels
1688 // =======================================================================
1689 Standard_Boolean OpenGl_Workspace::RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView,
1690 const Standard_ShortReal theOrigins[16],
1691 const Standard_ShortReal theDirects[16],
1692 const Standard_Integer theSizeX,
1693 const Standard_Integer theSizeY)
1695 if (myRaytraceRenderKernel == NULL || myComputeQueue == NULL)
1696 return Standard_False;
1698 ////////////////////////////////////////////////////////////////////////
1699 // Set kernel arguments
1701 cl_uint anIndex = 0;
1704 cl_int aLightSourceBufferSize = (cl_int)myRaytraceGeometry.Sources.size();
1706 anError |= clSetKernelArg (
1707 myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theSizeX);
1708 anError |= clSetKernelArg (
1709 myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theSizeY);
1710 anError |= clSetKernelArg (
1711 myRaytraceRenderKernel, anIndex++, sizeof(cl_float16), theOrigins);
1712 anError |= clSetKernelArg (
1713 myRaytraceRenderKernel, anIndex++, sizeof(cl_float16), theDirects);
1714 anError |= clSetKernelArg (
1715 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceEnvironment);
1716 anError |= clSetKernelArg (
1717 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImage);
1718 anError |= clSetKernelArg (
1719 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneNodeInfoBuffer);
1720 anError |= clSetKernelArg (
1721 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneMinPointBuffer);
1722 anError |= clSetKernelArg (
1723 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneMaxPointBuffer);
1724 anError |= clSetKernelArg (
1725 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectNodeInfoBuffer);
1726 anError |= clSetKernelArg (
1727 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectMinPointBuffer);
1728 anError |= clSetKernelArg (
1729 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectMaxPointBuffer);
1730 anError |= clSetKernelArg (
1731 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryTriangBuffer);
1732 anError |= clSetKernelArg (
1733 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryVertexBuffer);
1734 anError |= clSetKernelArg (
1735 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryNormalBuffer);
1736 anError |= clSetKernelArg (
1737 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceLightSourceBuffer);
1738 anError |= clSetKernelArg (
1739 myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceMaterialBuffer);
1740 anError |= clSetKernelArg (
1741 myRaytraceRenderKernel, anIndex++, sizeof(cl_float4), &myRaytraceGeometry.GlobalAmbient);
1742 anError |= clSetKernelArg (
1743 myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &aLightSourceBufferSize);
1744 anError |= clSetKernelArg (
1745 myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theCView.IsShadowsEnabled);
1746 anError |= clSetKernelArg (
1747 myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theCView.IsReflectionsEnabled);
1748 anError |= clSetKernelArg (
1749 myRaytraceRenderKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneEpsilon);
1750 anError |= clSetKernelArg (
1751 myRaytraceRenderKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneRadius);
1753 if (anError != CL_SUCCESS)
1755 const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of ray-tracing kernel!";
1757 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1758 GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
1760 return Standard_False;
1763 // Second-pass 'smoothing' kernel runs only if anti-aliasing is enabled
1764 if (theCView.IsAntialiasingEnabled)
1768 anError |= clSetKernelArg (
1769 myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theSizeX);
1770 anError |= clSetKernelArg (
1771 myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theSizeY);
1772 anError |= clSetKernelArg (
1773 myRaytraceSmoothKernel, anIndex++, sizeof(cl_float16), theOrigins);
1774 anError |= clSetKernelArg (
1775 myRaytraceSmoothKernel, anIndex++, sizeof(cl_float16), theDirects);
1776 anError |= clSetKernelArg (
1777 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImage);
1778 anError |= clSetKernelArg (
1779 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceEnvironment);
1780 anError |= clSetKernelArg (
1781 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImageAA);
1782 anError |= clSetKernelArg (
1783 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneNodeInfoBuffer);
1784 anError |= clSetKernelArg (
1785 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneMinPointBuffer);
1786 anError |= clSetKernelArg (
1787 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneMaxPointBuffer);
1788 anError |= clSetKernelArg (
1789 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectNodeInfoBuffer);
1790 anError |= clSetKernelArg (
1791 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectMinPointBuffer);
1792 anError |= clSetKernelArg (
1793 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectMaxPointBuffer);
1794 anError |= clSetKernelArg (
1795 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryTriangBuffer);
1796 anError |= clSetKernelArg (
1797 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryVertexBuffer);
1798 anError |= clSetKernelArg (
1799 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryNormalBuffer);
1800 anError |= clSetKernelArg (
1801 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceLightSourceBuffer);
1802 anError |= clSetKernelArg (
1803 myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceMaterialBuffer);
1804 anError |= clSetKernelArg (
1805 myRaytraceSmoothKernel, anIndex++, sizeof(cl_float4), &myRaytraceGeometry.GlobalAmbient);
1806 anError |= clSetKernelArg (
1807 myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &aLightSourceBufferSize);
1808 anError |= clSetKernelArg (
1809 myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theCView.IsShadowsEnabled);
1810 anError |= clSetKernelArg (
1811 myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theCView.IsReflectionsEnabled);
1812 anError |= clSetKernelArg (
1813 myRaytraceSmoothKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneEpsilon);
1814 anError |= clSetKernelArg (
1815 myRaytraceSmoothKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneRadius);
1817 if (anError != CL_SUCCESS)
1819 const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of smoothing kernel!";
1821 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1822 GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
1824 return Standard_False;
1828 ////////////////////////////////////////////////////////////////////////
1831 size_t aLocWorkSize[] = { myIsAmdComputePlatform ? 2 : 4, 32 };
1833 #ifdef OPENCL_GROUP_SIZE_TEST
1834 for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1)
1835 for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1)
1837 aLocWorkSize[0] = aLocX;
1838 aLocWorkSize[1] = aLocY;
1841 size_t aWorkSizeX = theSizeX;
1842 if (aWorkSizeX % aLocWorkSize[0] != 0)
1843 aWorkSizeX += aLocWorkSize[0] - aWorkSizeX % aLocWorkSize[0];
1845 size_t aWokrSizeY = theSizeY;
1846 if (aWokrSizeY % aLocWorkSize[1] != 0 )
1847 aWokrSizeY += aLocWorkSize[1] - aWokrSizeY % aLocWorkSize[1];
1849 size_t aTotWorkSize[] = { aWorkSizeX, aWokrSizeY };
1851 cl_event anEvent = NULL, anEventSmooth = NULL;
1853 anError = clEnqueueNDRangeKernel (myComputeQueue,
1854 myRaytraceRenderKernel, 2, NULL, aTotWorkSize, aLocWorkSize, 0, NULL, &anEvent);
1856 if (anError != CL_SUCCESS)
1858 const TCollection_ExtendedString aMessage = "Error! Failed to execute the ray-tracing kernel!";
1860 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1861 GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
1863 return Standard_False;
1866 clWaitForEvents (1, &anEvent);
1868 if (theCView.IsAntialiasingEnabled)
1870 size_t aLocWorkSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4,
1871 myIsAmdComputePlatform ? 8 : 32 };
1873 #ifdef OPENCL_GROUP_SIZE_TEST
1874 aLocWorkSizeSmooth[0] = aLocX;
1875 aLocWorkSizeSmooth[1] = aLocY;
1878 aWorkSizeX = theSizeX;
1879 if (aWorkSizeX % aLocWorkSizeSmooth[0] != 0)
1880 aWorkSizeX += aLocWorkSizeSmooth[0] - aWorkSizeX % aLocWorkSizeSmooth[0];
1882 size_t aWokrSizeY = theSizeY;
1883 if (aWokrSizeY % aLocWorkSizeSmooth[1] != 0 )
1884 aWokrSizeY += aLocWorkSizeSmooth[1] - aWokrSizeY % aLocWorkSizeSmooth[1];
1886 size_t aTotWorkSizeSmooth [] = { aWorkSizeX, aWokrSizeY };
1888 anError = clEnqueueNDRangeKernel (myComputeQueue, myRaytraceSmoothKernel,
1889 2, NULL, aTotWorkSizeSmooth, aLocWorkSizeSmooth, 0, NULL, &anEventSmooth);
1891 clWaitForEvents (1, &anEventSmooth);
1893 if (anError != CL_SUCCESS)
1895 const TCollection_ExtendedString aMessage = "Error! Failed to execute the smoothing kernel";
1897 myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
1898 GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
1900 return Standard_False;
1904 #if defined (RAY_TRACE_PRINT_INFO) || defined (OPENCL_GROUP_SIZE_TEST)
1906 static cl_ulong ttt1 = 10000000000;
1907 static cl_ulong ttt2 = 10000000000;
1909 cl_ulong aBegTime = 0;
1910 cl_ulong aEndTime = 0;
1912 clGetEventProfilingInfo (anEvent,
1913 CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL);
1914 clGetEventProfilingInfo (anEvent,
1915 CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL);
1917 ttt1 = aEndTime - aBegTime < ttt1 ? aEndTime - aBegTime : ttt1;
1919 std::cout << "\tRender time (ms): " << ttt1 / 1e6f << std::endl;
1921 if (theCView.IsAntialiasingEnabled)
1923 clGetEventProfilingInfo (anEventSmooth,
1924 CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL);
1925 clGetEventProfilingInfo (anEventSmooth,
1926 CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL);
1928 ttt2 = aEndTime - aBegTime < ttt2 ? aEndTime - aBegTime : ttt2;
1930 std::cout << "\tSmooth time (ms): " << ttt2 / 1e6f << std::endl;
1935 if (anEvent != NULL)
1936 clReleaseEvent (anEvent);
1938 if (anEventSmooth != NULL)
1939 clReleaseEvent (anEventSmooth);
1941 #ifdef OPENCL_GROUP_SIZE_TEST
1945 return Standard_True;
1948 // =======================================================================
1949 // function : ComputeInverseMatrix
1950 // purpose : Computes inversion of 4x4 floating-point matrix
1951 // =======================================================================
1952 template <typename T>
1953 void ComputeInverseMatrix (const T m[16], T inv[16])
1955 inv[ 0] = m[ 5] * (m[10] * m[15] - m[11] * m[14]) -
1956 m[ 9] * (m[ 6] * m[15] - m[ 7] * m[14]) -
1957 m[13] * (m[ 7] * m[10] - m[ 6] * m[11]);
1959 inv[ 1] = m[ 1] * (m[11] * m[14] - m[10] * m[15]) -
1960 m[ 9] * (m[ 3] * m[14] - m[ 2] * m[15]) -
1961 m[13] * (m[ 2] * m[11] - m[ 3] * m[10]);
1963 inv[ 2] = m[ 1] * (m[ 6] * m[15] - m[ 7] * m[14]) -
1964 m[ 5] * (m[ 2] * m[15] - m[ 3] * m[14]) -
1965 m[13] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]);
1967 inv[ 3] = m[ 1] * (m[ 7] * m[10] - m[ 6] * m[11]) -
1968 m[ 5] * (m[ 3] * m[10] - m[ 2] * m[11]) -
1969 m[ 9] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]);
1971 inv[ 4] = m[ 4] * (m[11] * m[14] - m[10] * m[15]) -
1972 m[ 8] * (m[ 7] * m[14] - m[ 6] * m[15]) -
1973 m[12] * (m[ 6] * m[11] - m[ 7] * m[10]);
1975 inv[ 5] = m[ 0] * (m[10] * m[15] - m[11] * m[14]) -
1976 m[ 8] * (m[ 2] * m[15] - m[ 3] * m[14]) -
1977 m[12] * (m[ 3] * m[10] - m[ 2] * m[11]);
1979 inv[ 6] = m[ 0] * (m[ 7] * m[14] - m[ 6] * m[15]) -
1980 m[ 4] * (m[ 3] * m[14] - m[ 2] * m[15]) -
1981 m[12] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]);
1983 inv[ 7] = m[ 0] * (m[ 6] * m[11] - m[ 7] * m[10]) -
1984 m[ 4] * (m[ 2] * m[11] - m[ 3] * m[10]) -
1985 m[ 8] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]);
1987 inv[ 8] = m[ 4] * (m[ 9] * m[15] - m[11] * m[13]) -
1988 m[ 8] * (m[ 5] * m[15] - m[ 7] * m[13]) -
1989 m[12] * (m[ 7] * m[ 9] - m[ 5] * m[11]);
1991 inv[ 9] = m[ 0] * (m[11] * m[13] - m[ 9] * m[15]) -
1992 m[ 8] * (m[ 3] * m[13] - m[ 1] * m[15]) -
1993 m[12] * (m[ 1] * m[11] - m[ 3] * m[ 9]);
1995 inv[10] = m[ 0] * (m[ 5] * m[15] - m[ 7] * m[13]) -
1996 m[ 4] * (m[ 1] * m[15] - m[ 3] * m[13]) -
1997 m[12] * (m[ 3] * m[ 5] - m[ 1] * m[ 7]);
1999 inv[11] = m[ 0] * (m[ 7] * m[ 9] - m[ 5] * m[11]) -
2000 m[ 4] * (m[ 3] * m[ 9] - m[ 1] * m[11]) -
2001 m[ 8] * (m[ 1] * m[ 7] - m[ 3] * m[ 5]);
2003 inv[12] = m[ 4] * (m[10] * m[13] - m[ 9] * m[14]) -
2004 m[ 8] * (m[ 6] * m[13] - m[ 5] * m[14]) -
2005 m[12] * (m[ 5] * m[10] - m[ 6] * m[ 9]);
2007 inv[13] = m[ 0] * (m[ 9] * m[14] - m[10] * m[13]) -
2008 m[ 8] * (m[ 1] * m[14] - m[ 2] * m[13]) -
2009 m[12] * (m[ 2] * m[ 9] - m[ 1] * m[10]);
2011 inv[14] = m[ 0] * (m[ 6] * m[13] - m[ 5] * m[14]) -
2012 m[ 4] * (m[ 2] * m[13] - m[ 1] * m[14]) -
2013 m[12] * (m[ 1] * m[ 6] - m[ 2] * m[ 5]);
2015 inv[15] = m[ 0] * (m[ 5] * m[10] - m[ 6] * m[ 9]) -
2016 m[ 4] * (m[ 1] * m[10] - m[ 2] * m[ 9]) -
2017 m[ 8] * (m[ 2] * m[ 5] - m[ 1] * m[ 6]);
2019 T det = m[0] * inv[ 0] +
2024 if (det == T (0.0)) return;
2026 det = T (1.0) / det;
2028 for (Standard_Integer i = 0; i < 16; ++i)
2032 // =======================================================================
2033 // function : GenerateCornerRays
2034 // purpose : Generates primary rays for corners of screen quad
2035 // =======================================================================
2036 void GenerateCornerRays (const GLdouble theInvModelProj[16],
2037 cl_float theOrigins[16],
2038 cl_float theDirects[16])
2040 Standard_Integer aOriginIndex = 0;
2041 Standard_Integer aDirectIndex = 0;
2043 for (Standard_Integer y = -1; y <= 1; y += 2)
2045 for (Standard_Integer x = -1; x <= 1; x += 2)
2047 BVH_Vec4f aOrigin (float(x),
2052 aOrigin = MatVecMult (theInvModelProj, aOrigin);
2053 aOrigin.x() = aOrigin.x() / aOrigin.w();
2054 aOrigin.y() = aOrigin.y() / aOrigin.w();
2055 aOrigin.z() = aOrigin.z() / aOrigin.w();
2058 BVH_Vec4f aDirect (float(x),
2063 aDirect = MatVecMult (theInvModelProj, aDirect);
2064 aDirect.x() = aDirect.x() / aDirect.w();
2065 aDirect.y() = aDirect.y() / aDirect.w();
2066 aDirect.z() = aDirect.z() / aDirect.w();
2069 aDirect = aDirect - aOrigin;
2071 GLdouble aInvLen = 1.f / sqrt (aDirect.x() * aDirect.x() +
2072 aDirect.y() * aDirect.y() +
2073 aDirect.z() * aDirect.z());
2075 theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.x());
2076 theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.y());
2077 theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.z());
2078 theOrigins [aOriginIndex++] = 1.f;
2080 theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.x() * aInvLen);
2081 theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.y() * aInvLen);
2082 theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.z() * aInvLen);
2083 theDirects [aDirectIndex++] = 0.f;
2088 // =======================================================================
2089 // function : Raytrace
2090 // purpose : Redraws the window using OpenCL ray tracing
2091 // =======================================================================
2092 Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView,
2093 const Standard_Integer theSizeX,
2094 const Standard_Integer theSizeY,
2095 const Tint theToSwap)
2098 return Standard_False;
2100 if (!ResizeRaytraceOutputBuffer (theSizeX, theSizeY))
2101 return Standard_False;
2103 if (!UpdateRaytraceEnvironmentMap())
2104 return Standard_False;
2106 if (!UpdateRaytraceGeometry (Standard_True))
2107 return Standard_False;
2109 // Get model-view and projection matrices
2110 TColStd_Array2OfReal theOrientation (0, 3, 0, 3);
2111 TColStd_Array2OfReal theViewMapping (0, 3, 0, 3);
2113 myView->GetMatrices (theOrientation, theViewMapping);
2115 GLdouble aOrientationMatrix[16];
2116 GLdouble aViewMappingMatrix[16];
2117 GLdouble aOrientationInvers[16];
2119 for (Standard_Integer j = 0; j < 4; ++j)
2120 for (Standard_Integer i = 0; i < 4; ++i)
2122 aOrientationMatrix [4 * j + i] = theOrientation (i, j);
2123 aViewMappingMatrix [4 * j + i] = theViewMapping (i, j);
2126 ComputeInverseMatrix (aOrientationMatrix, aOrientationInvers);
2128 if (!UpdateRaytraceLightSources (aOrientationInvers))
2129 return Standard_False;
2131 // Generate primary rays for corners of the screen quad
2132 glMatrixMode (GL_MODELVIEW);
2134 glLoadMatrixd (aViewMappingMatrix);
2135 glMultMatrixd (aOrientationMatrix);
2137 GLdouble aModelProject[16];
2138 GLdouble aInvModelProj[16];
2140 glGetDoublev (GL_MODELVIEW_MATRIX, aModelProject);
2142 ComputeInverseMatrix (aModelProject, aInvModelProj);
2144 GLfloat aOrigins[16];
2145 GLfloat aDirects[16];
2147 GenerateCornerRays (aInvModelProj,
2151 // Compute ray-traced image using OpenCL kernel
2152 cl_mem anImages[] = { myRaytraceOutputImage, myRaytraceOutputImageAA };
2153 cl_int anError = clEnqueueAcquireGLObjects (myComputeQueue,
2156 clFinish (myComputeQueue);
2158 if (myIsRaytraceDataValid)
2160 RunRaytraceOpenCLKernels (theCView,
2167 anError |= clEnqueueReleaseGLObjects (myComputeQueue,
2170 clFinish (myComputeQueue);
2173 glPushAttrib (GL_ENABLE_BIT |
2175 GL_COLOR_BUFFER_BIT |
2176 GL_DEPTH_BUFFER_BIT);
2178 glDisable (GL_DEPTH_TEST);
2180 if (NamedStatus & OPENGL_NS_WHITEBACK)
2182 glClearColor (1.0f, 1.0f, 1.0f, 1.0f);
2186 glClearColor (myBgColor.rgb[0],
2192 glClear (GL_COLOR_BUFFER_BIT);
2194 Handle(OpenGl_Workspace) aWorkspace (this);
2195 myView->DrawBackground (aWorkspace);
2197 // Draw dummy quad to show result image
2198 glEnable (GL_COLOR_MATERIAL);
2199 glEnable (GL_BLEND);
2201 glDisable (GL_DEPTH_TEST);
2203 glBlendFunc (GL_ONE, GL_SRC_ALPHA);
2205 glEnable (GL_TEXTURE_RECTANGLE);
2207 glMatrixMode (GL_PROJECTION);
2210 glMatrixMode (GL_MODELVIEW);
2213 glColor3f (1.0f, 1.0f, 1.0f);
2215 if (!theCView.IsAntialiasingEnabled)
2216 myRaytraceOutputTexture->Bind (myGlContext);
2218 myRaytraceOutputTextureAA->Bind (myGlContext);
2220 if (myIsRaytraceDataValid)
2224 glTexCoord2i ( 0, 0); glVertex2f (-1.f, -1.f);
2225 glTexCoord2i ( 0, theSizeY); glVertex2f (-1.f, 1.f);
2226 glTexCoord2i (theSizeX, theSizeY); glVertex2f ( 1.f, 1.f);
2227 glTexCoord2i (theSizeX, 0); glVertex2f ( 1.f, -1.f);
2237 GetGlContext()->SwapBuffers();
2238 myBackBufferRestored = Standard_False;
2243 return Standard_True;