diff --git a/samples/tcl/raytrace.tcl b/samples/tcl/raytrace.tcl index b07e238edc..efe00e42e2 100644 --- a/samples/tcl/raytrace.tcl +++ b/samples/tcl/raytrace.tcl @@ -30,10 +30,8 @@ vsetcolorbg 255 255 255 vfit # set ray tracing -if { [regexp {HAVE_OPENCL} [dversion]] } { - puts "Trying raytrace mode..." - if { ! [catch {vraytrace 1}] } { - vtextureenv on 1 - vsetraytracemode shad=1 refl=1 aa=1 - } +puts "Trying raytrace mode..." +if { ! [catch {vraytrace 1}] } { + vtextureenv on 1 + vsetraytracemode shad=1 refl=1 aa=1 } diff --git a/src/BVH/BVH_BinnedBuilder.lxx b/src/BVH/BVH_BinnedBuilder.lxx index 85511a93cb..8a705784c3 100644 --- a/src/BVH/BVH_BinnedBuilder.lxx +++ b/src/BVH/BVH_BinnedBuilder.lxx @@ -282,6 +282,10 @@ void BVH_BinnedBuilder::BuildNode (BVH_Set* theSet, || theBVH->Level (aChildIndex) >= BVH_Builder::myMaxTreeDepth; if (!isLeaf) + { BVH_Builder::myTasksQueue.Append (aChildIndex); + } + + BVH_Builder::UpdateDepth (theBVH, theBVH->Level (aChildIndex)); } } diff --git a/src/BVH/BVH_Builder.hxx b/src/BVH/BVH_Builder.hxx index b5e3181831..09093bf390 100644 --- a/src/BVH/BVH_Builder.hxx +++ b/src/BVH/BVH_Builder.hxx @@ -54,6 +54,16 @@ protected: BVH_Tree* theBVH, const Standard_Integer theTask); + //! Updates depth of constructed BVH tree. + void UpdateDepth (BVH_Tree* theBVH, + const Standard_Integer theLevel) + { + if (theLevel > theBVH->myDepth) + { + theBVH->myDepth = theLevel; + } + } + protected: Standard_Integer myMaxTreeDepth; //!< Maximum depth of constructed BVH diff --git a/src/BVH/BVH_SweepPlaneBuilder.lxx b/src/BVH/BVH_SweepPlaneBuilder.lxx index 7e8fcf68a3..62ef763b68 100644 --- a/src/BVH/BVH_SweepPlaneBuilder.lxx +++ b/src/BVH/BVH_SweepPlaneBuilder.lxx @@ -184,5 +184,7 @@ void BVH_SweepPlaneBuilder::BuildNode (BVH_Set* theSet, { BVH_Builder::myTasksQueue.Append (aChildIndex); } + + BVH_Builder::UpdateDepth (theBVH, theBVH->Level (aChildIndex)); } } diff --git a/src/BVH/BVH_Tree.hxx b/src/BVH/BVH_Tree.hxx index 7b7b661cd5..2b309d40d1 100644 --- a/src/BVH/BVH_Tree.hxx +++ b/src/BVH/BVH_Tree.hxx @@ -20,6 +20,8 @@ #include +template class BVH_Builder; + //! Stores parameters of bounding volume hierarchy (BVH). //! Bounding volume hierarchy (BVH) organizes geometric objects in //! the tree based on spatial relationships. Each node in the tree @@ -31,12 +33,20 @@ template class BVH_Tree { + friend class BVH_Builder; + public: typedef typename BVH_Box::BVH_VecNt BVH_VecNt; public: + //! Creates new empty BVH tree. + BVH_Tree() : myDepth (0) + { + // + } + //! Returns minimum point of the given node. BVH_VecNt& MinPoint (const Standard_Integer theNodeIndex) { @@ -151,6 +161,12 @@ public: return BVHTools::ArrayOp::Size (myNodeInfoBuffer); } + //! Returns depth of BVH tree from last build. + Standard_Integer Depth() const + { + return myDepth; + } + public: //! Removes all BVH nodes. @@ -232,6 +248,9 @@ protected: //! Array of node data records. BVH_Array4i myNodeInfoBuffer; + //! Depth of constructed tree. + Standard_Integer myDepth; + }; #include diff --git a/src/BVH/BVH_Tree.lxx b/src/BVH/BVH_Tree.lxx index 35aaaeb75a..fe9e2a3a80 100644 --- a/src/BVH/BVH_Tree.lxx +++ b/src/BVH/BVH_Tree.lxx @@ -20,6 +20,8 @@ template void BVH_Tree::Clear() { + myDepth = 0; + BVHTools::ArrayOp::Clear (myMinPointBuffer); BVHTools::ArrayOp::Clear (myMaxPointBuffer); diff --git a/src/Graphic3d/Graphic3d_CView.hxx b/src/Graphic3d/Graphic3d_CView.hxx index 3c49890934..e5cf2998af 100644 --- a/src/Graphic3d/Graphic3d_CView.hxx +++ b/src/Graphic3d/Graphic3d_CView.hxx @@ -100,7 +100,7 @@ public: WasRedrawnGL (0), IsRaytracing (0), IsShadowsEnabled (1), - IsReflectionsEnabled (1), + IsReflectionsEnabled (0), IsAntialiasingEnabled (0) { memset(&DefWindow,0,sizeof(DefWindow)); diff --git a/src/OpenGl/EXTERNLIB b/src/OpenGl/EXTERNLIB index a586f17029..21580fc276 100755 --- a/src/OpenGl/EXTERNLIB +++ b/src/OpenGl/EXTERNLIB @@ -6,7 +6,6 @@ CSF_objc CSF_Appkit CSF_IOKit CSF_OpenGlLibs -CSF_OPENCL CSF_AviLibs CSF_FREETYPE CSF_GL2PS diff --git a/src/OpenGl/FILES b/src/OpenGl/FILES index 04ff3873af..5619791097 100755 --- a/src/OpenGl/FILES +++ b/src/OpenGl/FILES @@ -140,10 +140,8 @@ OpenGl_ShaderStates.cxx Handle_OpenGl_ShaderObject.hxx Handle_OpenGl_ShaderProgram.hxx Handle_OpenGl_ShaderManager.hxx -OpenGl_Cl.hxx OpenGl_SceneGeometry.hxx OpenGl_SceneGeometry.cxx -OpenGl_RaytraceSource.cxx OpenGl_Workspace_Raytrace.cxx OpenGl_Flipper.hxx OpenGl_Flipper.cxx diff --git a/src/OpenGl/OpenGl_Cl.hxx b/src/OpenGl/OpenGl_Cl.hxx deleted file mode 100755 index 2524e61609..0000000000 --- a/src/OpenGl/OpenGl_Cl.hxx +++ /dev/null @@ -1,29 +0,0 @@ -// Created on: 2013-10-15 -// Created by: Denis BOGOLEPOV -// Copyright (c) 2013-2014 OPEN CASCADE SAS -// -// This file is part of Open CASCADE Technology software library. -// -// This library is free software; you can redistribute it and/or modify it under -// the terms of the GNU Lesser General Public License version 2.1 as published -// by the Free Software Foundation, with special exception defined in the file -// OCCT_LGPL_EXCEPTION.txt. Consult the file LICENSE_LGPL_21.txt included in OCCT -// distribution for complete text of the license and disclaimer of any warranty. -// -// Alternatively, this file may be used under the terms of Open CASCADE -// commercial license or contractual agreement. - -#ifndef _OpenGl_Cl_H__ -#define _OpenGl_Cl_H__ - -// cl_gl.h includes OpenGL headers - make sure our stuff is included in right order -#include - -#if defined(__APPLE__) || defined(__MACOSX) - #include -#else - #include - #include -#endif - -#endif // _OpenGl_Cl_H__ diff --git a/src/OpenGl/OpenGl_GraphicDriver.cxx b/src/OpenGl/OpenGl_GraphicDriver.cxx index 3052c0d78b..879e61f0aa 100755 --- a/src/OpenGl/OpenGl_GraphicDriver.cxx +++ b/src/OpenGl/OpenGl_GraphicDriver.cxx @@ -171,34 +171,6 @@ Standard_Boolean OpenGl_GraphicDriver::SetImmediateModeDrawToFront (const Graphi return Standard_False; } -// ======================================================================= -// function : GetOpenClDeviceInfo -// purpose : Returns information about device used for computations -// ======================================================================= -#ifndef HAVE_OPENCL - -Standard_Boolean OpenGl_GraphicDriver::GetOpenClDeviceInfo (const Graphic3d_CView&, - NCollection_DataMap&) -{ - return Standard_False; -} - -#else - -Standard_Boolean OpenGl_GraphicDriver::GetOpenClDeviceInfo (const Graphic3d_CView& theCView, - NCollection_DataMap& theInfo) -{ - - if (theCView.ViewId == -1 || theCView.ptrView == NULL) - { - return Standard_False; - } - - return reinterpret_cast (theCView.ptrView)->WS->GetOpenClDeviceInfo (theInfo); -} - -#endif - // ======================================================================= // function : DisplayImmediateStructure // purpose : diff --git a/src/OpenGl/OpenGl_GraphicDriver.hxx b/src/OpenGl/OpenGl_GraphicDriver.hxx index 92caa80979..8492e551e4 100644 --- a/src/OpenGl/OpenGl_GraphicDriver.hxx +++ b/src/OpenGl/OpenGl_GraphicDriver.hxx @@ -308,10 +308,6 @@ public: Standard_EXPORT OpenGl_UserDrawCallback_t& UserDrawCallback(); public: - - //! Returns information about OpenCL device used for computations. - Standard_EXPORT Standard_Boolean GetOpenClDeviceInfo (const Graphic3d_CView& theCView, - NCollection_DataMap& theInfo); //! Method to retrieve valid GL context. //! Could return NULL-handle if no window created by this driver. diff --git a/src/OpenGl/OpenGl_Group.cxx b/src/OpenGl/OpenGl_Group.cxx index 8d23f9a5d9..10af74d528 100644 --- a/src/OpenGl/OpenGl_Group.cxx +++ b/src/OpenGl/OpenGl_Group.cxx @@ -13,10 +13,6 @@ // Alternatively, this file may be used under the terms of Open CASCADE // commercial license or contractual agreement. -#ifdef HAVE_CONFIG_H - #include -#endif - #include #include @@ -118,7 +114,6 @@ void OpenGl_Group::UpdateAspectFace (const Standard_Boolean theIsGlobal) AddElement (anAspectFace); } -#ifdef HAVE_OPENCL if (myIsRaytracable) { ++myModificationState; @@ -128,7 +123,6 @@ void OpenGl_Group::UpdateAspectFace (const Standard_Boolean theIsGlobal) aStruct->UpdateStateWithAncestorStructures(); } } -#endif } // ======================================================================= @@ -300,7 +294,6 @@ void OpenGl_Group::AddElement (OpenGl_Element* theElem) (myLast? myLast->next : myFirst) = aNode; myLast = aNode; -#ifdef HAVE_OPENCL if (OpenGl_Raytrace::IsRaytracedElement (aNode)) { myModificationState++; @@ -313,7 +306,6 @@ void OpenGl_Group::AddElement (OpenGl_Element* theElem) aStruct->SetRaytracableWithAncestorStructures(); } } -#endif } // ======================================================================= diff --git a/src/OpenGl/OpenGl_LayerList.cxx b/src/OpenGl/OpenGl_LayerList.cxx index b80d9d6b1b..f494bf19e9 100644 --- a/src/OpenGl/OpenGl_LayerList.cxx +++ b/src/OpenGl/OpenGl_LayerList.cxx @@ -13,10 +13,6 @@ // Alternatively, this file may be used under the terms of Open CASCADE // commercial license or contractual agreement. -#ifdef HAVE_CONFIG_H - #include -#endif - #include #include @@ -176,12 +172,10 @@ void OpenGl_LayerList::RemoveStructure (const OpenGl_Structure *theStructure, { myNbStructures--; -#ifdef HAVE_OPENCL if (theStructure->IsRaytracable()) { myModificationState++; } -#endif return; } @@ -199,12 +193,10 @@ void OpenGl_LayerList::RemoveStructure (const OpenGl_Structure *theStructure, { myNbStructures--; -#ifdef HAVE_OPENCL if (theStructure->IsRaytracable()) { myModificationState++; } -#endif return; } diff --git a/src/OpenGl/OpenGl_LayerList.hxx b/src/OpenGl/OpenGl_LayerList.hxx index 32424495d2..848337d305 100644 --- a/src/OpenGl/OpenGl_LayerList.hxx +++ b/src/OpenGl/OpenGl_LayerList.hxx @@ -77,14 +77,10 @@ class OpenGl_LayerList //! Returns the set of OpenGL Z-layers. const OpenGl_SequenceOfLayers& Layers() const { return myLayers; } - -#ifdef HAVE_OPENCL //! Returns structure modification state (for ray-tracing). Standard_Size ModificationState() const { return myModificationState; } -#endif - private: //! Get default layer @@ -98,9 +94,7 @@ class OpenGl_LayerList Standard_Integer myNbPriorities; Standard_Integer myNbStructures; -#ifdef HAVE_OPENCL mutable Standard_Size myModificationState; -#endif public: DEFINE_STANDARD_ALLOC diff --git a/src/OpenGl/OpenGl_RaytraceSource.cxx b/src/OpenGl/OpenGl_RaytraceSource.cxx deleted file mode 100755 index 654ed2e229..0000000000 --- a/src/OpenGl/OpenGl_RaytraceSource.cxx +++ /dev/null @@ -1,1163 +0,0 @@ -// Created on: 2013-10-16 -// Created by: Denis BOGOLEPOV -// Copyright (c) 2013-2014 OPEN CASCADE SAS -// -// This file is part of Open CASCADE Technology software library. -// -// This library is free software; you can redistribute it and/or modify it under -// the terms of the GNU Lesser General Public License version 2.1 as published -// by the Free Software Foundation, with special exception defined in the file -// OCCT_LGPL_EXCEPTION.txt. Consult the file LICENSE_LGPL_21.txt included in OCCT -// distribution for complete text of the license and disclaimer of any warranty. -// -// Alternatively, this file may be used under the terms of Open CASCADE -// commercial license or contractual agreement. - -#ifdef HAVE_CONFIG_H - #include -#endif - -#ifdef HAVE_OPENCL - -#define EOL "\n" - -extern const char THE_RAY_TRACE_OPENCL_SOURCE[] = - - ///////////////////////////////////////////////////////////////////////////////////////// - // Specific data types - EOL - //! Stores ray parameters. - EOL" typedef struct __SRay" - EOL" {" - EOL" float4 Origin;" - EOL" float4 Direct;" - EOL" }" - EOL" SRay;" - EOL - //! Stores parameters of intersection point. - EOL" typedef struct __SIntersect" - EOL" {" - EOL" float4 Normal;" - EOL" float Time;" - EOL" float U;" - EOL" float V;" - EOL" }" - EOL" SIntersect;" - EOL - EOL - ///////////////////////////////////////////////////////////////////////////////////////// - // Some useful vector constants - EOL - EOL" #define ZERO ( float4 )( 0.f, 0.f, 0.f, 0.f )" - EOL" #define UNIT ( float4 )( 1.f, 1.f, 1.f, 0.f )" - EOL - EOL" #define AXIS_X ( float4 )( 1.f, 0.f, 0.f, 0.f )" - EOL" #define AXIS_Y ( float4 )( 0.f, 1.f, 0.f, 0.f )" - EOL" #define AXIS_Z ( float4 )( 0.f, 0.f, 1.f, 0.f )" - EOL - ///////////////////////////////////////////////////////////////////////////////////////// - // Support functions - EOL - // ======================================================================= - // function : GenerateRay - // purpose : Generates primary ray for current work item - // ======================================================================= - EOL" void GenerateRay (SRay* theRay," - EOL" const float theX," - EOL" const float theY," - EOL" const int theSizeX," - EOL" const int theSizeY," - EOL" const float16 theOrigins," - EOL" const float16 theDirects)" - EOL" {" - EOL" float2 aPixel = (float2) (theX / (float)theSizeX," - EOL" theY / (float)theSizeY);" - EOL - EOL" float4 aP0 = mix (theOrigins.lo.lo, theOrigins.lo.hi, aPixel.x);" - EOL" float4 aP1 = mix (theOrigins.hi.lo, theOrigins.hi.hi, aPixel.x);" - EOL - EOL" theRay->Origin = mix (aP0, aP1, aPixel.y);" - EOL - EOL" aP0 = mix (theDirects.lo.lo, theDirects.lo.hi, aPixel.x);" - EOL" aP1 = mix (theDirects.hi.lo, theDirects.hi.hi, aPixel.x);" - EOL - EOL" theRay->Direct = mix (aP0, aP1, aPixel.y);" - EOL" }" - EOL - EOL - ///////////////////////////////////////////////////////////////////////////////////////// - // Functions for compute ray-object intersection - EOL - EOL" #define _OOEPS_ exp2 (-80.0f)" - EOL - // ======================================================================= - // function : IntersectSphere - // purpose : Computes ray-sphere intersection - // ======================================================================= - EOL" float IntersectSphere (const SRay* theRay, float theRadius)" - EOL" {" - EOL" float aDdotD = dot (theRay->Direct.xyz, theRay->Direct.xyz);" - EOL" float aDdotO = dot (theRay->Direct.xyz, theRay->Origin.xyz);" - EOL" float aOdotO = dot (theRay->Origin.xyz, theRay->Origin.xyz);" - EOL - EOL" float aD = aDdotO * aDdotO - aDdotD * (aOdotO - theRadius * theRadius);" - EOL - EOL" if (aD > 0.f)" - EOL" {" - EOL" float aTime = (-aDdotO + native_sqrt (aD)) * (1.f / aDdotD);" - EOL - EOL" return aTime > 0.f ? aTime : MAXFLOAT;" - EOL" }" - EOL - EOL" return MAXFLOAT;" - EOL" }" - EOL - // ======================================================================= - // function : IntersectBox - // purpose : Computes ray-box intersection (slab test) - // ======================================================================= - EOL" float IntersectBox (const SRay* theRay," - EOL" float4 theMinPoint," - EOL" float4 theMaxPoint)" - EOL" {" - EOL" const float4 aInvDirect = (float4)(" - EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?" - EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x))," - EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?" - EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y))," - EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?" - EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z))," - EOL" 0.f);" - EOL - EOL" const float4 aTime0 = (theMinPoint - theRay->Origin) * aInvDirect;" - EOL" const float4 aTime1 = (theMaxPoint - theRay->Origin) * aInvDirect;" - EOL - EOL" const float4 aTimeMax = max (aTime0, aTime1);" - EOL" const float4 aTimeMin = min (aTime0, aTime1);" - EOL - EOL" const float theTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" - EOL" const float theTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" - EOL - EOL" return (theTimeStart <= theTimeFinal) && (theTimeFinal >= 0.f) ? theTimeStart : MAXFLOAT;" - EOL" }" - EOL - // ======================================================================= - // function : IntersectNodes - // purpose : Computes intersection of ray with two child nodes (boxes) - // ======================================================================= - EOL" void IntersectNodes (const SRay* theRay," - EOL" float4 theMinPoint0," - EOL" float4 theMaxPoint0," - EOL" float4 theMinPoint1," - EOL" float4 theMaxPoint1," - EOL" float* theTimeStart0," - EOL" float* theTimeStart1," - EOL" float theMaxTime)" - EOL" {" - EOL" const float4 aInvDirect = (float4)(" - EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?" - EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x))," - EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?" - EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y))," - EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?" - EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z))," - EOL" 0.f);" - EOL - EOL" float4 aTime0 = (theMinPoint0 - theRay->Origin) * aInvDirect;" - EOL" float4 aTime1 = (theMaxPoint0 - theRay->Origin) * aInvDirect;" - EOL - EOL" float4 aTimeMax = max (aTime0, aTime1);" - EOL" float4 aTimeMin = min (aTime0, aTime1);" - EOL - EOL" aTime0 = (theMinPoint1 - theRay->Origin) * aInvDirect;" - EOL" aTime1 = (theMaxPoint1 - theRay->Origin) * aInvDirect;" - EOL - EOL" float aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" - EOL" float aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" - EOL - EOL" aTimeMax = max (aTime0, aTime1);" - EOL" aTimeMin = min (aTime0, aTime1);" - EOL - EOL" *theTimeStart0 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)" - EOL" ? aTimeStart : -MAXFLOAT;" - EOL - EOL" aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" - EOL" aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" - EOL - EOL" *theTimeStart1 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)" - EOL" ? aTimeStart : -MAXFLOAT;" - EOL" }" - EOL - // ======================================================================= - // function : IntersectTriangle - // purpose : Computes ray-triangle intersection (branchless version) - // ======================================================================= - EOL" float IntersectTriangle (const SRay* theRay," - EOL" const float4 thePoint0," - EOL" const float4 thePoint1," - EOL" const float4 thePoint2," - EOL" float4* theNormal," - EOL" float* theU," - EOL" float* theV)" - EOL" {" - EOL" const float4 aEdge0 = thePoint1 - thePoint0;" - EOL" const float4 aEdge1 = thePoint0 - thePoint2;" - EOL - EOL" *theNormal = cross (aEdge1, aEdge0);" - EOL - EOL" const float4 aEdge2 = (1.f / dot (*theNormal, theRay->Direct)) * (thePoint0 - theRay->Origin);" - EOL - EOL" const float aTime = dot (*theNormal, aEdge2);" - EOL - EOL" const float4 theVec = cross (theRay->Direct, aEdge2);" - EOL - EOL" *theU = dot (theVec, aEdge1);" - EOL" *theV = dot (theVec, aEdge0);" - EOL - EOL" return (aTime >= 0.f) & (*theU >= 0.f) & (*theV >= 0.f) & (*theU + *theV <= 1.f) ? aTime : MAXFLOAT;" - EOL" }" - EOL - ///////////////////////////////////////////////////////////////////////////////////////// - // Support shading functions - EOL - EOL" const sampler_t EnvironmentSampler =" - EOL" CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_REPEAT | CLK_FILTER_LINEAR;" - EOL - // ======================================================================= - // function : SmoothNormal - // purpose : Interpolates normal across the triangle - // ======================================================================= - EOL" float4 SmoothNormal (__global float4* theNormals," - EOL" const SIntersect* theHit," - EOL" const int4 theIndices)" - EOL" {" - EOL" float4 aNormal0 = theNormals[theIndices.x]," - EOL" aNormal1 = theNormals[theIndices.y]," - EOL" aNormal2 = theNormals[theIndices.z];" - EOL - EOL" return fast_normalize (aNormal1 * theHit->U +" - EOL" aNormal2 * theHit->V +" - EOL" aNormal0 * (1.f - theHit->U - theHit->V));" - EOL" }" - EOL - // ======================================================================= - // function : Shade - // purpose : Computes Phong-based illumination - // ======================================================================= - EOL" float4 Shade (const float4 theMatDiff," - EOL" const float4 theMatSpec," - EOL" const float4 theLight," - EOL" const float4 theView," - EOL" const float4 theNormal," - EOL" const float4 theIntens," - EOL" const float theTranspr)" - EOL" {" - EOL" float aLambert = dot (theNormal, theLight);" - EOL - EOL" aLambert = theTranspr > 0.f ? fabs (aLambert) : aLambert;" - EOL - EOL" if (aLambert > 0.f)" - EOL" {" - EOL" const float4 aReflect = 2.f * dot (theLight, theNormal) * theNormal - theLight;" - EOL - EOL" const float aSpecular = pow (max (dot (aReflect.xyz, theView.xyz), 0.f), theMatSpec.w);" - EOL - EOL" return theIntens * (theMatDiff * aLambert + theMatSpec * aSpecular);" - EOL" }" - EOL - EOL" return ZERO;" - EOL" }" - EOL - // ======================================================================= - // function : Latlong - // purpose : Converts world direction to environment texture coordinates - // ======================================================================= - EOL" float2 Latlong (const float4 thePoint, const float theRadius)" - EOL" {" - EOL" float aPsi = acospi (-thePoint.y / theRadius);" - EOL" float aPhi = atan2pi (thePoint.z, thePoint.x);" - EOL - EOL" aPhi = (aPhi < 0.f) ? aPhi + 2.f : aPhi;" - EOL - EOL" return (float2) (aPhi * 0.5f, aPsi);" - EOL" }" - EOL - ///////////////////////////////////////////////////////////////////////////////////////// - // Core ray tracing function - EOL - // ======================================================================= - // function : push - // purpose : Pushes BVH node index to local stack - // ======================================================================= - EOL" void push (uint* theStack, char* thePos, const uint theValue)" - EOL" {" - EOL" (*thePos)++;" - EOL" theStack[*thePos] = theValue;" - EOL" }" - EOL - // ======================================================================= - // function : pop - // purpose : Pops BVH node index from local stack - // ======================================================================= - EOL" void pop (uint* theStack, char* thePos, uint* theValue)" - EOL" {" - EOL" *theValue = theStack[*thePos];" - EOL" (*thePos)--;" - EOL" }" - EOL - // ======================================================================= - // function : ObjectNearestHit - // purpose : Finds intersection with nearest object triangle - // ======================================================================= - EOL" int4 ObjectNearestHit (const SRay* theRay," - EOL" SIntersect* theIntersect," - EOL" __global int4* theObjectNodeInfoBuffer," - EOL" __global float4* theObjectMinPointBuffer," - EOL" __global float4* theObjectMaxPointBuffer," - EOL" __global int4* theGeometryTriangBuffer," - EOL" __global float4* theGeometryVertexBuffer)" - EOL" {" - EOL" uint aStack [32];" - EOL - EOL" char aHead = -1;" // stack pointer - EOL" uint aNode = 0;" // node to visit - EOL - EOL" const float4 aInvDirect = (float4) (" - EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?" - EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x))," - EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?" - EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y))," - EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?" - EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z))," - EOL" 0.f);" - EOL - EOL" int4 aTriangleIndex = (int4) (-1);" - EOL - EOL" float aTimeExit;" - EOL" float aTimeMin1;" - EOL" float aTimeMin2;" - EOL - EOL" while (true)" - EOL" {" - EOL" const int3 aData = theObjectNodeInfoBuffer[aNode].xyz;" - EOL - EOL" if (aData.x == 0)" // if inner node - EOL" {" - EOL" float4 aNodeMin = theObjectMinPointBuffer[aData.y];" - EOL" float4 aNodeMax = theObjectMaxPointBuffer[aData.y];" - EOL - EOL" float4 aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;" - EOL" float4 aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;" - EOL - EOL" float4 aTimeMax = max (aTime0, aTime1);" - EOL" float4 aTimeMin = min (aTime0, aTime1);" - EOL - EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" - EOL" aTimeMin1 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" - EOL - EOL" const bool aHitLft = (aTimeMin1 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin1 <= theIntersect->Time);" - EOL - EOL" aNodeMin = theObjectMinPointBuffer[aData.z];" - EOL" aNodeMax = theObjectMaxPointBuffer[aData.z];" - EOL - EOL" aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;" - EOL" aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;" - EOL - EOL" aTimeMax = max (aTime0, aTime1);" - EOL" aTimeMin = min (aTime0, aTime1);" - EOL - EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" - EOL" aTimeMin2 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" - EOL - EOL" const bool aHitRgh = (aTimeMin2 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin2 <= theIntersect->Time);" - EOL - EOL" if (aHitLft & aHitRgh)" - EOL" {" - EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;" - EOL - EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);" - EOL" }" - EOL" else" - EOL" {" - EOL" if (aHitLft | aHitRgh)" - EOL" {" - EOL" aNode = aHitLft ? aData.y : aData.z;" - EOL" }" - EOL" else" - EOL" {" - EOL" if (aHead < 0)" - EOL" return aTriangleIndex;" - EOL - EOL" pop (aStack, &aHead, &aNode);" - EOL" }" - EOL" }" - EOL" }" - EOL" else " // if leaf node - EOL" {" - EOL" for (int aTriIdx = aData.y; aTriIdx <= aData.z; ++aTriIdx)" - EOL" {" - EOL" const int4 aTestTriangle = theGeometryTriangBuffer[aTriIdx];" - EOL - EOL" const float4 aPoint0 = theGeometryVertexBuffer[aTestTriangle.x];" - EOL" const float4 aPoint1 = theGeometryVertexBuffer[aTestTriangle.y];" - EOL" const float4 aPoint2 = theGeometryVertexBuffer[aTestTriangle.z];" - EOL - EOL" float4 aNormal; float aU, aV;" - EOL - EOL" float aTime = IntersectTriangle (theRay," - EOL" aPoint0," - EOL" aPoint1," - EOL" aPoint2," - EOL" &aNormal," - EOL" &aU," - EOL" &aV);" - EOL - EOL" if (aTime < theIntersect->Time)" - EOL" {" - EOL" aTriangleIndex = aTestTriangle;" - EOL" theIntersect->Normal = aNormal;" - EOL" theIntersect->Time = aTime;" - EOL" theIntersect->U = aU;" - EOL" theIntersect->V = aV;" - EOL" }" - EOL" }" - EOL - EOL" if (aHead < 0)" - EOL" return aTriangleIndex;" - EOL - EOL" pop (aStack, &aHead, &aNode);" - EOL" }" - EOL" }" - EOL - EOL" return aTriangleIndex;" - EOL" }" - EOL - // ======================================================================= - // function : ObjectAnyHit - // purpose : Finds intersection with any object triangle - // ======================================================================= - EOL" float ObjectAnyHit (const SRay* theRay," - EOL" __global int4* theObjectNodeInfoBuffer," - EOL" __global float4* theObjectMinPointBuffer," - EOL" __global float4* theObjectMaxPointBuffer," - EOL" __global int4* theGeometryTriangBuffer," - EOL" __global float4* theGeometryVertexBuffer," - EOL" const float theDistance)" - EOL" {" - EOL" uint aStack [32];" - EOL - EOL" char aHead = -1;" // stack pointer - EOL" uint aNode = 0;" // node to visit - EOL - EOL" const float4 aInvDirect = (float4) (" - EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?" - EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x))," - EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?" - EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y))," - EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?" - EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z))," - EOL" 0.f);" - EOL - EOL" float aTimeExit;" - EOL" float aTimeMin1;" - EOL" float aTimeMin2;" - EOL - EOL" while (true)" - EOL" {" - EOL" const int3 aData = theObjectNodeInfoBuffer[aNode].xyz;" - EOL - EOL" if (aData.x == 0)" // if inner node - EOL" {" - EOL" float4 aNodeMin = theObjectMinPointBuffer[aData.y];" - EOL" float4 aNodeMax = theObjectMaxPointBuffer[aData.y];" - EOL - EOL" float4 aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;" - EOL" float4 aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;" - EOL - EOL" float4 aTimeMax = max (aTime0, aTime1);" - EOL" float4 aTimeMin = min (aTime0, aTime1);" - EOL - EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" - EOL" aTimeMin1 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" - EOL - EOL" const bool aHitLft = (aTimeMin1 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin1 <= theDistance);" - EOL - EOL" aNodeMin = theObjectMinPointBuffer[aData.z];" - EOL" aNodeMax = theObjectMaxPointBuffer[aData.z];" - EOL - EOL" aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;" - EOL" aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;" - EOL - EOL" aTimeMax = max (aTime0, aTime1);" - EOL" aTimeMin = min (aTime0, aTime1);" - EOL - EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));" - EOL" aTimeMin2 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));" - EOL - EOL" const bool aHitRgh = (aTimeMin2 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin2 <= theDistance);" - EOL - EOL" if (aHitLft & aHitRgh)" - EOL" {" - EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;" - EOL - EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);" - EOL" }" - EOL" else" - EOL" {" - EOL" if (aHitLft | aHitRgh)" - EOL" {" - EOL" aNode = aHitLft ? aData.y : aData.z;" - EOL" }" - EOL" else" - EOL" {" - EOL" if (aHead < 0)" - EOL" return 1.f;" - EOL - EOL" pop (aStack, &aHead, &aNode);" - EOL" }" - EOL" }" - EOL" }" - EOL" else " // if leaf node - EOL" {" - EOL" for (int aTriIdx = aData.y; aTriIdx <= aData.z; ++aTriIdx)" - EOL" {" - EOL" const int4 aTestTriangle = theGeometryTriangBuffer[aTriIdx];" - EOL - EOL" const float4 aPoint0 = theGeometryVertexBuffer[aTestTriangle.x];" - EOL" const float4 aPoint1 = theGeometryVertexBuffer[aTestTriangle.y];" - EOL" const float4 aPoint2 = theGeometryVertexBuffer[aTestTriangle.z];" - EOL - EOL" float4 aNormal; float aU, aV;" - EOL - EOL" float aTime = IntersectTriangle (theRay," - EOL" aPoint0," - EOL" aPoint1," - EOL" aPoint2," - EOL" &aNormal," - EOL" &aU," - EOL" &aV);" - EOL - EOL" if (aTime < theDistance)" - EOL" {" - EOL" return 0.f;" - EOL" }" - EOL" }" - EOL - EOL" if (aHead < 0)" - EOL" return 1.f;" - EOL - EOL" pop (aStack, &aHead, &aNode);" - EOL" }" - EOL" }" - EOL - EOL" return 1.f;" - EOL" }" - EOL - // ======================================================================= - // function : NearestHit - // purpose : Finds intersection with nearest scene triangle - // ======================================================================= - EOL" int4 NearestHit (const SRay* theRay," - EOL" SIntersect* theIntersect," - EOL" __global int4* theSceneNodeInfoBuffer," - EOL" __global float4* theSceneMinPointBuffer," - EOL" __global float4* theSceneMaxPointBuffer," - EOL" __global int4* theObjectNodeInfoBuffer," - EOL" __global float4* theObjectMinPointBuffer," - EOL" __global float4* theObjectMaxPointBuffer," - EOL" __global int4* theGeometryTriangBuffer," - EOL" __global float4* theGeometryVertexBuffer)" - EOL" {" - EOL" theIntersect->Time = MAXFLOAT;" - EOL - EOL" uint aStack [16];" - EOL - EOL" char aHead = -1;" // stack pointer - EOL" uint aNode = 0;" // node to visit - EOL - EOL" int4 aNearestTriangle = (int4) (-1);" - EOL - EOL" while (true)" - EOL" {" - EOL" const int4 aData = theSceneNodeInfoBuffer[aNode];" - EOL - EOL" if (aData.x != 0)" // if leaf node - EOL" {" - EOL" const float4 aNodeMin = theSceneMinPointBuffer[aNode];" - EOL" const float4 aNodeMax = theSceneMaxPointBuffer[aNode];" - EOL - EOL" if (IntersectBox (theRay, aNodeMin, aNodeMax) <= theIntersect->Time)" - EOL" {" - EOL" int4 anIndex = ObjectNearestHit (theRay," - EOL" theIntersect," - EOL" theObjectNodeInfoBuffer + aData.y," - EOL" theObjectMinPointBuffer + aData.y," - EOL" theObjectMaxPointBuffer + aData.y," - EOL" theGeometryTriangBuffer + aData.w," - EOL" theGeometryVertexBuffer + aData.z);" - EOL - EOL" if (anIndex.x != -1)" - EOL" aNearestTriangle = (int4) (anIndex.x + aData.z," - EOL" anIndex.y + aData.z," - EOL" anIndex.z + aData.z," - EOL" anIndex.w);" - EOL" }" - EOL - EOL" if (aHead < 0)" - EOL" return aNearestTriangle;" - EOL - EOL" pop (aStack, &aHead, &aNode);" - EOL" }" - EOL" else " // if inner node - EOL" {" - EOL" float4 aNodeMinLft = theSceneMinPointBuffer[aData.y];" - EOL" float4 aNodeMinRgh = theSceneMinPointBuffer[aData.z];" - EOL" float4 aNodeMaxLft = theSceneMaxPointBuffer[aData.y];" - EOL" float4 aNodeMaxRgh = theSceneMaxPointBuffer[aData.z];" - EOL - EOL" float aTimeMin1;" - EOL" float aTimeMin2;" - EOL - EOL" IntersectNodes (theRay," - EOL" aNodeMinLft," - EOL" aNodeMaxLft," - EOL" aNodeMinRgh," - EOL" aNodeMaxRgh," - EOL" &aTimeMin1," - EOL" &aTimeMin2," - EOL" theIntersect->Time);" - EOL - EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);" - EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);" - EOL - EOL" if (aHitLft & aHitRgh)" - EOL" {" - EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;" - EOL - EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);" - EOL" }" - EOL" else" - EOL" {" - EOL" if (aHitLft | aHitRgh)" - EOL" {" - EOL" aNode = aHitLft ? aData.y : aData.z;" - EOL" }" - EOL" else" - EOL" {" - EOL" if (aHead < 0)" - EOL" return aNearestTriangle;" - EOL - EOL" pop (aStack, &aHead, &aNode);" - EOL" }" - EOL" }" - EOL" }" - EOL" }" - EOL - EOL" return aNearestTriangle;" - EOL" }" - EOL - // ======================================================================= - // function : AnyHit - // purpose : Finds intersection with any scene triangle - // ======================================================================= - EOL" float AnyHit (const SRay* theRay," - EOL" __global int4* theSceneNodeInfoBuffer," - EOL" __global float4* theSceneMinPointBuffer," - EOL" __global float4* theSceneMaxPointBuffer," - EOL" __global int4* theObjectNodeInfoBuffer," - EOL" __global float4* theObjectMinPointBuffer," - EOL" __global float4* theObjectMaxPointBuffer," - EOL" __global int4* theGeometryTriangBuffer," - EOL" __global float4* theGeometryVertexBuffer," - EOL" const float theDistance)" - EOL" {" - EOL" uint aStack [16];" - EOL - EOL" char aHead = -1;" // stack pointer - EOL" uint aNode = 0;" // node to visit - EOL - EOL" while (true)" - EOL" {" - EOL" const int4 aData = theSceneNodeInfoBuffer[aNode];" - EOL - EOL" if (aData.x != 0)" // if leaf node - EOL" {" - EOL" const float4 aNodeMin = theSceneMinPointBuffer[aNode];" - EOL" const float4 aNodeMax = theSceneMaxPointBuffer[aNode];" - EOL - EOL" if (IntersectBox (theRay, aNodeMin, aNodeMax) <= theDistance)" - EOL" {" - EOL" if (0.f == ObjectAnyHit (theRay," - EOL" theObjectNodeInfoBuffer + aData.y," - EOL" theObjectMinPointBuffer + aData.y," - EOL" theObjectMaxPointBuffer + aData.y," - EOL" theGeometryTriangBuffer + aData.w," - EOL" theGeometryVertexBuffer + aData.z," - EOL" theDistance))" - EOL" {" - EOL" return 0.f;" - EOL" }" - EOL" }" - EOL - EOL" if (aHead < 0)" - EOL" return 1.f;" - EOL - EOL" pop (aStack, &aHead, &aNode);" - EOL" }" - EOL" else" // if inner node - EOL" {" - EOL" float4 aNodeMinLft = theSceneMinPointBuffer[aData.y];" - EOL" float4 aNodeMinRgh = theSceneMinPointBuffer[aData.z];" - EOL" float4 aNodeMaxLft = theSceneMaxPointBuffer[aData.y];" - EOL" float4 aNodeMaxRgh = theSceneMaxPointBuffer[aData.z];" - EOL - EOL" float aTimeMin1;" - EOL" float aTimeMin2;" - EOL - EOL" IntersectNodes (theRay," - EOL" aNodeMinLft," - EOL" aNodeMaxLft," - EOL" aNodeMinRgh," - EOL" aNodeMaxRgh," - EOL" &aTimeMin1," - EOL" &aTimeMin2," - EOL" theDistance);" - EOL - EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);" - EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);" - EOL - EOL" if (aHitLft & aHitRgh)" - EOL" {" - EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;" - EOL - EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);" - EOL" }" - EOL" else" - EOL" {" - EOL" if (aHitLft | aHitRgh)" - EOL" {" - EOL" aNode = aHitLft ? aData.y : aData.z;" - EOL" }" - EOL" else" - EOL" {" - EOL" if (aHead < 0)" - EOL" return 1.f;" - EOL - EOL" pop (aStack, &aHead, &aNode);" - EOL" }" - EOL" }" - EOL" }" - EOL" }" - EOL" }" - EOL - EOL" #define _MAX_DEPTH_ 5" - EOL - EOL" #define THRESHOLD (float4) (0.1f, 0.1f, 0.1f, 1.f)" - EOL - EOL" #define LIGHT_POS(Buffer, LightID) Buffer[2 * LightID + 1]" - EOL" #define LIGHT_RAD(Buffer, LightID) Buffer[2 * LightID + 0]" - EOL - EOL" #define MATERIAL_AMBN(Buffer, TriangleID) Buffer[7 * TriangleID.w + 0]" - EOL" #define MATERIAL_DIFF(Buffer, TriangleID) Buffer[7 * TriangleID.w + 1]" - EOL" #define MATERIAL_SPEC(Buffer, TriangleID) Buffer[7 * TriangleID.w + 2]" - EOL" #define MATERIAL_EMIS(Buffer, TriangleID) Buffer[7 * TriangleID.w + 3]" - EOL" #define MATERIAL_REFL(Buffer, TriangleID) Buffer[7 * TriangleID.w + 4]" - EOL" #define MATERIAL_REFR(Buffer, TriangleID) Buffer[7 * TriangleID.w + 5]" - EOL" #define MATERIAL_TRAN(Buffer, TriangleID) Buffer[7 * TriangleID.w + 6]" - EOL - // ======================================================================= - // function : Radiance - // purpose : Computes color of specified ray - // ======================================================================= - EOL" float4 Radiance (SRay* theRay," - EOL" __read_only image2d_t theEnvMap," - EOL" __global int4* theSceneNodeInfoBuffer," - EOL" __global float4* theSceneMinPointBuffer," - EOL" __global float4* theSceneMaxPointBuffer," - EOL" __global int4* theObjectNodeInfoBuffer," - EOL" __global float4* theObjectMinPointBuffer," - EOL" __global float4* theObjectMaxPointBuffer," - EOL" __global int4* theGeometryTriangBuffer," - EOL" __global float4* theGeometryVertexBuffer," - EOL" __global float4* theGeometryNormalBuffer," - EOL" __global float4* theLightSourceBuffer," - EOL" __global float4* theMaterialBuffer," - EOL" const float4 theGlobalAmbient," - EOL" const int theLightBufferSize," - EOL" const int theShadowsEnabled," - EOL" const int theReflectEnabled," - EOL" const float theSceneEpsilon," - EOL" const float theSceneRadius)" - EOL" {" - EOL" float4 aResult = (float4) (0.f, 0.f, 0.f, 0.f);" - EOL" float4 aWeight = (float4) (1.f, 1.f, 1.f, 1.f);" - EOL - EOL" SIntersect aHit;" - EOL - EOL" for (int aDepth = 0; aDepth < _MAX_DEPTH_; ++aDepth)" - EOL" {" - EOL" int4 aTriangle = NearestHit (theRay," - EOL" &aHit," - EOL" theSceneNodeInfoBuffer," - EOL" theSceneMinPointBuffer," - EOL" theSceneMaxPointBuffer," - EOL" theObjectNodeInfoBuffer," - EOL" theObjectMinPointBuffer," - EOL" theObjectMaxPointBuffer," - EOL" theGeometryTriangBuffer," - EOL" theGeometryVertexBuffer);" - EOL - EOL" if (aTriangle.x < 0.f)" - EOL" {" - EOL" if (aWeight.w != 0.f)" - EOL" break;" - EOL - EOL" float aTime = IntersectSphere (theRay, theSceneRadius);" - EOL - EOL" if (aTime != MAXFLOAT)" - EOL" {" - EOL" aResult += aWeight * read_imagef (theEnvMap, EnvironmentSampler," - EOL" Latlong (theRay->Origin + theRay->Direct * aTime, theSceneRadius));" - EOL" }" - EOL - EOL" return (float4) (aResult.x," - EOL" aResult.y," - EOL" aResult.z," - EOL" aWeight.w);" - EOL" }" - EOL - EOL // Compute geometric normal - EOL" float4 aGeomNormal = aHit.Normal; aGeomNormal = fast_normalize (aGeomNormal);" - EOL - EOL // Compute interpolated normal - EOL" float4 aNormal = SmoothNormal (theGeometryNormalBuffer, &aHit, aTriangle);" - EOL - EOL // Compute intersection point - EOL" float4 aPoint = theRay->Direct * aHit.Time + theRay->Origin;" - EOL - EOL" float4 aMaterAmb = MATERIAL_AMBN (theMaterialBuffer, aTriangle);" - EOL" float4 aMaterTrn = MATERIAL_TRAN (theMaterialBuffer, aTriangle);" - EOL - EOL" aResult += aWeight * theGlobalAmbient * aMaterAmb *" - EOL" (aMaterTrn.x * max (fabs (dot (theRay->Direct, aNormal)), 0.5f));" - EOL - EOL" for (int nLight = 0; nLight < theLightBufferSize; ++nLight)" - EOL" {" - EOL" float4 aLightPosition = LIGHT_POS (theLightSourceBuffer, nLight);" - EOL - EOL" SRay aShadow;" - EOL" aShadow.Direct = aLightPosition;" - EOL - EOL" float aLightDistance = MAXFLOAT;" - EOL" if (aLightPosition.w != 0.f)" - EOL" {" - EOL" aLightDistance = length (aLightPosition - aPoint);" - EOL" aShadow.Direct = (aLightPosition - aPoint) * (1.f / aLightDistance);" - EOL" }" - EOL - EOL" aShadow.Origin = aPoint + aShadow.Direct * theSceneEpsilon +" - EOL" aGeomNormal * copysign (theSceneEpsilon, dot (aGeomNormal, aShadow.Direct));" - EOL - EOL" float aVisibility = 1.f;" - EOL - EOL" if (theShadowsEnabled)" - EOL" {" - EOL" aVisibility = AnyHit (&aShadow," - EOL" theSceneNodeInfoBuffer," - EOL" theSceneMinPointBuffer," - EOL" theSceneMaxPointBuffer," - EOL" theObjectNodeInfoBuffer," - EOL" theObjectMinPointBuffer," - EOL" theObjectMaxPointBuffer," - EOL" theGeometryTriangBuffer," - EOL" theGeometryVertexBuffer," - EOL" aLightDistance);" - EOL" }" - EOL - EOL" if (aVisibility > 0.f)" - EOL" {" - EOL" aResult += aMaterTrn.x * aWeight * Shade (MATERIAL_DIFF (theMaterialBuffer, aTriangle)," - EOL" MATERIAL_SPEC (theMaterialBuffer, aTriangle)," - EOL" aShadow.Direct," - EOL" -theRay->Direct," - EOL" aNormal," - EOL" LIGHT_RAD (theLightSourceBuffer, nLight)," - EOL" aMaterTrn.y);" - EOL" }" - EOL" }" - EOL - EOL" if (aMaterTrn.y > 0.f)" - EOL" {" - EOL" aWeight *= aMaterTrn.y;" - EOL" }" - EOL" else" - EOL" {" - EOL" aWeight *= theReflectEnabled ? MATERIAL_REFL (theMaterialBuffer, aTriangle) : ZERO;" - EOL - EOL" float4 aDirect = theRay->Direct - 2.f * dot (theRay->Direct, aNormal) * aNormal;" - EOL - EOL" float aDdotN = dot (aDirect, aGeomNormal);" - EOL" if (aDdotN < 0.f)" - EOL" theRay->Direct -= 2.f * dot (theRay->Direct, aGeomNormal) * aGeomNormal;" - EOL" else" - EOL" theRay->Direct = aDirect;" - EOL" }" - EOL - EOL" if (all (islessequal (aWeight, THRESHOLD)))" - EOL" {" - EOL" return (float4) (aResult.x," - EOL" aResult.y," - EOL" aResult.z," - EOL" aWeight.w);" - EOL" }" - EOL - EOL" theRay->Origin = theRay->Direct * theSceneEpsilon + aPoint;" - EOL" }" - EOL - EOL" return (float4) (aResult.x," - EOL" aResult.y," - EOL" aResult.z," - EOL" aWeight.w);" - EOL" }" - EOL - /////////////////////////////////////////////////////////////////////////////// - // Ray tracing kernel functions - EOL - // ======================================================================= - // function : RaytraceRender - // purpose : Computes pixel color using ray-tracing - // ======================================================================= - EOL" __kernel void RaytraceRender (const int theSizeX," - EOL" const int theSizeY," - EOL" const float16 theOrigins," - EOL" const float16 theDirects," - EOL" __read_only image2d_t theEnvMap," - EOL" __write_only image2d_t theOutput," - EOL" __global int4* theSceneNodeInfoBuffer," - EOL" __global float4* theSceneMinPointBuffer," - EOL" __global float4* theSceneMaxPointBuffer," - EOL" __global int4* theObjectNodeInfoBuffer," - EOL" __global float4* theObjectMinPointBuffer," - EOL" __global float4* theObjectMaxPointBuffer," - EOL" __global int4* theGeometryTriangBuffer," - EOL" __global float4* theGeometryVertexBuffer," - EOL" __global float4* theGeometryNormalBuffer," - EOL" __global float4* theLightSourceBuffer," - EOL" __global float4* theMaterialBuffer," - EOL" const float4 theGlobalAmbient," - EOL" const int theLightBufferSize," - EOL" const int theShadowsEnabled," - EOL" const int theReflectEnabled," - EOL" const float theSceneEpsilon," - EOL" const float theSceneRadius)" - EOL" {" - EOL" const int aPixelX = get_global_id (0);" - EOL" const int aPixelY = get_global_id (1);" - EOL - EOL" if (aPixelX >= theSizeX || aPixelY >= theSizeY)" - EOL" return;" - EOL - EOL" private SRay aRay;" - EOL - EOL" GenerateRay (&aRay," - EOL" aPixelX," - EOL" aPixelY," - EOL" theSizeX," - EOL" theSizeY," - EOL" theOrigins," - EOL" theDirects);" - EOL - EOL" float4 aColor = (float4) (0.f, 0.f, 0.f, 1.f);" - EOL - EOL" float aTimeStart = IntersectBox (&aRay, theSceneMinPointBuffer[0], theSceneMaxPointBuffer[0]);" - EOL - EOL" if (aTimeStart != MAXFLOAT)" - EOL" {" - EOL" aRay.Origin += aRay.Direct * max (aTimeStart - theSceneEpsilon, 0.f);" - EOL - EOL" aColor = clamp (Radiance (&aRay," - EOL" theEnvMap," - EOL" theSceneNodeInfoBuffer," - EOL" theSceneMinPointBuffer," - EOL" theSceneMaxPointBuffer," - EOL" theObjectNodeInfoBuffer," - EOL" theObjectMinPointBuffer," - EOL" theObjectMaxPointBuffer," - EOL" theGeometryTriangBuffer," - EOL" theGeometryVertexBuffer," - EOL" theGeometryNormalBuffer," - EOL" theLightSourceBuffer," - EOL" theMaterialBuffer," - EOL" theGlobalAmbient," - EOL" theLightBufferSize," - EOL" theShadowsEnabled," - EOL" theReflectEnabled," - EOL" theSceneEpsilon," - EOL" theSceneRadius), 0.f, 1.f);" - EOL" }" - EOL - EOL" write_imagef (theOutput, (int2) (aPixelX, aPixelY), aColor);" - EOL" }" - EOL - EOL" const sampler_t OutputSampler =" - EOL" CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;" - EOL - EOL" #define _LUM_DELTA_ 0.085f" - EOL - EOL" #define AA_MAX 0.559017f" - EOL" #define AA_MIN 0.186339f" - EOL - // ======================================================================= - // function : RaytraceSmooth - // purpose : Performs adaptive sub-pixel rendering - // ======================================================================= - EOL" __kernel void RaytraceSmooth (const int theSizeX," - EOL" const int theSizeY," - EOL" const float16 theOrigins," - EOL" const float16 theDirects," - EOL" __read_only image2d_t theInput," - EOL" __read_only image2d_t theEnvMap," - EOL" __write_only image2d_t theOutput," - EOL" __global int4* theSceneNodeInfoBuffer," - EOL" __global float4* theSceneMinPointBuffer," - EOL" __global float4* theSceneMaxPointBuffer," - EOL" __global int4* theObjectNodeInfoBuffer," - EOL" __global float4* theObjectMinPointBuffer," - EOL" __global float4* theObjectMaxPointBuffer," - EOL" __global int4* theGeometryTriangBuffer," - EOL" __global float4* theGeometryVertexBuffer," - EOL" __global float4* theGeometryNormalBuffer," - EOL" __global float4* theLightSourceBuffer," - EOL" __global float4* theMaterialBuffer," - EOL" const float4 theGlobalAmbient," - EOL" const int theLightBufferSize," - EOL" const int theShadowsEnabled," - EOL" const int theReflectEnabled," - EOL" const float theSceneEpsilon," - EOL" const float theSceneRadius)" - EOL" {" - EOL" const int aPixelX = get_global_id (0);" - EOL" const int aPixelY = get_global_id (1);" - EOL - EOL" if (aPixelX >= theSizeX || aPixelY >= theSizeY)" - EOL" return;" - EOL - EOL" float4 aClr0 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 0, aPixelY + 0));" - EOL" float4 aClr1 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 0, aPixelY - 1));" - EOL" float4 aClr2 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 0, aPixelY + 1));" - EOL - EOL" float4 aClr3 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 1, aPixelY + 0));" - EOL" float4 aClr4 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 1, aPixelY - 1));" - EOL" float4 aClr5 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 1, aPixelY + 1));" - EOL - EOL" float4 aClr6 = read_imagef (theInput, OutputSampler, (float2) (aPixelX - 1, aPixelY + 0));" - EOL" float4 aClr7 = read_imagef (theInput, OutputSampler, (float2) (aPixelX - 1, aPixelY - 1));" - EOL" float4 aClr8 = read_imagef (theInput, OutputSampler, (float2) (aPixelX - 1, aPixelY + 1));" - EOL - EOL" bool render = fabs (aClr1.w - aClr0.w) > _LUM_DELTA_ ||" - EOL" fabs (aClr2.w - aClr0.w) > _LUM_DELTA_ ||" - EOL" fabs (aClr3.w - aClr0.w) > _LUM_DELTA_ ||" - EOL" fabs (aClr4.w - aClr0.w) > _LUM_DELTA_ ||" - EOL" fabs (aClr5.w - aClr0.w) > _LUM_DELTA_ ||" - EOL" fabs (aClr6.w - aClr0.w) > _LUM_DELTA_ ||" - EOL" fabs (aClr7.w - aClr0.w) > _LUM_DELTA_ ||" - EOL" fabs (aClr8.w - aClr0.w) > _LUM_DELTA_;" - EOL - EOL" if (!render)" - EOL" {" - EOL" aClr1 = (aClr1.w == 1.f) ? -UNIT : aClr1;" - EOL" aClr2 = (aClr2.w == 1.f) ? -UNIT : aClr2;" - EOL" aClr3 = (aClr3.w == 1.f) ? -UNIT : aClr3;" - EOL" aClr4 = (aClr4.w == 1.f) ? -UNIT : aClr4;" - EOL" aClr5 = (aClr5.w == 1.f) ? -UNIT : aClr5;" - EOL" aClr6 = (aClr6.w == 1.f) ? -UNIT : aClr6;" - EOL" aClr7 = (aClr7.w == 1.f) ? -UNIT : aClr7;" - EOL" aClr8 = (aClr8.w == 1.f) ? -UNIT : aClr8;" - EOL - EOL" float aLum = (aClr0.w == 1.f) ? -1.f : (0.2126f * aClr0.x + 0.7152f * aClr0.y + 0.0722f * aClr0.z);" - EOL - EOL" render = fabs (0.2126f * aClr1.x + 0.7152f * aClr1.y + 0.0722f * aClr1.z - aLum) > _LUM_DELTA_ ||" - EOL" fabs (0.2126f * aClr2.x + 0.7152f * aClr2.y + 0.0722f * aClr2.z - aLum) > _LUM_DELTA_ ||" - EOL" fabs (0.2126f * aClr3.x + 0.7152f * aClr3.y + 0.0722f * aClr3.z - aLum) > _LUM_DELTA_ ||" - EOL" fabs (0.2126f * aClr4.x + 0.7152f * aClr4.y + 0.0722f * aClr4.z - aLum) > _LUM_DELTA_ ||" - EOL" fabs (0.2126f * aClr5.x + 0.7152f * aClr5.y + 0.0722f * aClr5.z - aLum) > _LUM_DELTA_ ||" - EOL" fabs (0.2126f * aClr6.x + 0.7152f * aClr6.y + 0.0722f * aClr6.z - aLum) > _LUM_DELTA_ ||" - EOL" fabs (0.2126f * aClr7.x + 0.7152f * aClr7.y + 0.0722f * aClr7.z - aLum) > _LUM_DELTA_ ||" - EOL" fabs (0.2126f * aClr8.x + 0.7152f * aClr8.y + 0.0722f * aClr8.z - aLum) > _LUM_DELTA_;" - EOL" }" - EOL - EOL" float4 aColor = clamp (aClr0, 0.f, 1.f);" - EOL - EOL" private SRay aRay;" - EOL - EOL" const float4 aBoxMin = theSceneMinPointBuffer[0];" - EOL" const float4 aBoxMax = theSceneMaxPointBuffer[0];" - EOL - EOL" if (render)" - EOL" {" - EOL" for (int aSample = 0; aSample <= 3; ++aSample)" - EOL" {" - EOL" float fX = aPixelX, fY = aPixelY;" - EOL - EOL" if (aSample == 0)" - EOL" {" - EOL" fX -= AA_MIN; fY -= AA_MAX;" - EOL" }" - EOL" else if (aSample == 1)" - EOL" {" - EOL" fX -= AA_MAX; fY += AA_MIN;" - EOL" }" - EOL" else if (aSample == 2)" - EOL" {" - EOL" fX += AA_MIN; fY += AA_MAX;" - EOL" }" - EOL" else" - EOL" {" - EOL" fX += AA_MAX; fY -= AA_MIN;" - EOL" }" - EOL - EOL" GenerateRay (&aRay," - EOL" fX," - EOL" fY," - EOL" theSizeX," - EOL" theSizeY," - EOL" theOrigins," - EOL" theDirects);" - EOL - EOL" float aTimeStart = IntersectBox (&aRay, aBoxMin, aBoxMax);" - EOL - EOL" if (aTimeStart != MAXFLOAT)" - EOL" {" - EOL" aRay.Origin += aRay.Direct * max (aTimeStart - theSceneEpsilon, 0.f);" - EOL - EOL" aColor += clamp (Radiance (&aRay," - EOL" theEnvMap," - EOL" theSceneNodeInfoBuffer," - EOL" theSceneMinPointBuffer," - EOL" theSceneMaxPointBuffer," - EOL" theObjectNodeInfoBuffer," - EOL" theObjectMinPointBuffer," - EOL" theObjectMaxPointBuffer," - EOL" theGeometryTriangBuffer," - EOL" theGeometryVertexBuffer," - EOL" theGeometryNormalBuffer," - EOL" theLightSourceBuffer," - EOL" theMaterialBuffer," - EOL" theGlobalAmbient," - EOL" theLightBufferSize," - EOL" theShadowsEnabled," - EOL" theReflectEnabled," - EOL" theSceneEpsilon," - EOL" theSceneRadius), 0.f, 1.f);" - EOL" }" - EOL" else" - EOL" aColor += (float4) (0.f, 0.f, 0.f, 1.f);" - EOL" }" - EOL - EOL" aColor *= 1.f / 5.f;" - EOL" }" - EOL - EOL" write_imagef (theOutput, (int2) (aPixelX, aPixelY), aColor);" - EOL" }"; - -#endif diff --git a/src/OpenGl/OpenGl_SceneGeometry.cxx b/src/OpenGl/OpenGl_SceneGeometry.cxx index b0bd4ab93d..73ecc43403 100755 --- a/src/OpenGl/OpenGl_SceneGeometry.cxx +++ b/src/OpenGl/OpenGl_SceneGeometry.cxx @@ -13,12 +13,6 @@ // Alternatively, this file may be used under the terms of Open CASCADE // commercial license or contractual agreement. -#ifdef HAVE_CONFIG_H - #include -#endif - -#ifdef HAVE_OPENCL - #include #ifdef HAVE_TBB @@ -206,6 +200,8 @@ Standard_Boolean OpenGl_RaytraceGeometry::ProcessAcceleration() OpenGL_BVHParallelBuilder (this)); #endif + myBottomLevelTreeDepth = 0; + for (Standard_Integer anObjectIdx = 0; anObjectIdx < Size(); ++anObjectIdx) { OpenGl_TriangleSet* aTriangleSet = dynamic_cast ( @@ -216,6 +212,8 @@ Standard_Boolean OpenGl_RaytraceGeometry::ProcessAcceleration() Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(), "Error! Failed to update bottom-level BVH of OpenGL element", Standard_False); + + myBottomLevelTreeDepth = Max (myBottomLevelTreeDepth, aTriangleSet->BVH()->Depth()); } #ifdef BVH_PRINT_INFO @@ -242,6 +240,8 @@ Standard_Boolean OpenGl_RaytraceGeometry::ProcessAcceleration() Standard_ASSERT_RETURN (!aBVH.IsNull(), "Error! Failed to update high-level BVH of ray-tracing scene", Standard_False); + myHighLevelTreeDepth = aBVH->Depth(); + Standard_Integer aVerticesOffset = 0; Standard_Integer aElementsOffset = 0; Standard_Integer aBVHNodesOffset = 0; @@ -387,5 +387,3 @@ namespace OpenGl_Raytrace return Standard_False; } } - -#endif diff --git a/src/OpenGl/OpenGl_SceneGeometry.hxx b/src/OpenGl/OpenGl_SceneGeometry.hxx index 36330fac70..6a6f793fcc 100755 --- a/src/OpenGl/OpenGl_SceneGeometry.hxx +++ b/src/OpenGl/OpenGl_SceneGeometry.hxx @@ -16,8 +16,6 @@ #ifndef _OpenGl_SceneGeometry_Header #define _OpenGl_SceneGeometry_Header -#ifdef HAVE_OPENCL - #include #include #include @@ -124,13 +122,13 @@ class OpenGl_TriangleSet : public BVH_Triangulation { public: - //! Array of vertex normals. - BVH_Array4f Normals; + BVH_Array4f Normals; //!< Array of vertex normals public: //! Creates new OpenGL element triangulation. OpenGl_TriangleSet() + : BVH_Triangulation() { // } @@ -161,12 +159,15 @@ public: NCollection_StdAllocator > Materials; //! Global ambient from all light sources. - BVH_Vec4f GlobalAmbient; + BVH_Vec4f Ambient; public: //! Creates uninitialized ray-tracing geometry. OpenGl_RaytraceGeometry() + : BVH_Geometry(), + myHighLevelTreeDepth (0), + myBottomLevelTreeDepth (0) { // } @@ -204,7 +205,24 @@ public: //! If the node index is not valid the function returns NULL. //! @note Can be used after processing acceleration structure. OpenGl_TriangleSet* TriangleSet (Standard_Integer theNodeIdx); + + //! Returns depth of high-level scene BVH from last build. + Standard_Integer HighLevelTreeDepth() const + { + return myHighLevelTreeDepth; + } + + //! Returns maximum depth of bottom-level scene BVHs from last build. + Standard_Integer BottomLevelTreeDepth() const + { + return myBottomLevelTreeDepth; + } + +protected: + + Standard_Integer myHighLevelTreeDepth; //!< Depth of high-level scene BVH from last build + Standard_Integer myBottomLevelTreeDepth; //!< Maximum depth of bottom-level scene BVHs from last build + }; #endif -#endif diff --git a/src/OpenGl/OpenGl_ShaderProgram.cxx b/src/OpenGl/OpenGl_ShaderProgram.cxx index 9a97a6bbf4..a4366f6087 100755 --- a/src/OpenGl/OpenGl_ShaderProgram.cxx +++ b/src/OpenGl/OpenGl_ShaderProgram.cxx @@ -374,17 +374,19 @@ Standard_Boolean OpenGl_ShaderProgram::Link (const Handle(OpenGl_Context)& theCt return Standard_False; } - theCtx->core20->glLinkProgram (myProgramID); - GLint aStatus = GL_FALSE; + theCtx->core20->glLinkProgram (myProgramID); theCtx->core20->glGetProgramiv (myProgramID, GL_LINK_STATUS, &aStatus); + if (aStatus == GL_FALSE) + { + return Standard_False; + } for (GLint aVar = 0; aVar < OpenGl_OCCT_NUMBER_OF_STATE_VARIABLES; ++aVar) { myStateLocations[aVar] = GetUniformLocation (theCtx, PredefinedKeywords[aVar]); } - - return aStatus != GL_FALSE; + return Standard_True; } // ======================================================================= @@ -644,6 +646,130 @@ Standard_Boolean OpenGl_ShaderProgram::GetAttribute (const Handle(OpenGl_Context return Standard_True; } +// ======================================================================= +// function : SetAttributeName +// purpose : +// ======================================================================= +Standard_Boolean OpenGl_ShaderProgram::SetAttributeName (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + const GLchar* theName) +{ + theCtx->core20fwd->glBindAttribLocation (myProgramID, theIndex, theName); + return Standard_True; +} + +// ======================================================================= +// function : SetAttribute +// purpose : +// ======================================================================= +Standard_Boolean OpenGl_ShaderProgram::SetAttribute (const Handle(OpenGl_Context)& theCtx, + const GLchar* theName, + GLfloat theValue) +{ + return SetAttribute (theCtx, GetAttributeLocation (theCtx, theName), theValue); +} + +// ======================================================================= +// function : SetAttribute +// purpose : +// ======================================================================= +Standard_Boolean OpenGl_ShaderProgram::SetAttribute (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + GLfloat theValue) +{ + if (myProgramID == NO_PROGRAM || theIndex == INVALID_LOCATION) + { + return Standard_False; + } + + theCtx->core20fwd->glVertexAttrib1f (theIndex, theValue); + return Standard_True; +} + +// ======================================================================= +// function : SetAttribute +// purpose : +// ======================================================================= +Standard_Boolean OpenGl_ShaderProgram::SetAttribute (const Handle(OpenGl_Context)& theCtx, + const GLchar* theName, + const OpenGl_Vec2& theValue) +{ + return SetAttribute (theCtx, GetAttributeLocation (theCtx, theName), theValue); +} + +// ======================================================================= +// function : SetAttribute +// purpose : +// ======================================================================= +Standard_Boolean OpenGl_ShaderProgram::SetAttribute (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + const OpenGl_Vec2& theValue) +{ + if (myProgramID == NO_PROGRAM || theIndex == INVALID_LOCATION) + { + return Standard_False; + } + + theCtx->core20fwd->glVertexAttrib2fv (theIndex, theValue); + return Standard_True; +} + +// ======================================================================= +// function : SetAttribute +// purpose : +// ======================================================================= +Standard_Boolean OpenGl_ShaderProgram::SetAttribute (const Handle(OpenGl_Context)& theCtx, + const GLchar* theName, + const OpenGl_Vec3& theValue) +{ + return SetAttribute (theCtx, GetAttributeLocation (theCtx, theName), theValue); +} + +// ======================================================================= +// function : SetAttribute +// purpose : +// ======================================================================= +Standard_Boolean OpenGl_ShaderProgram::SetAttribute (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + const OpenGl_Vec3& theValue) +{ + if (myProgramID == NO_PROGRAM || theIndex == INVALID_LOCATION) + { + return Standard_False; + } + + theCtx->core20fwd->glVertexAttrib3fv (theIndex, theValue); + return Standard_True; +} + +// ======================================================================= +// function : SetAttribute +// purpose : +// ======================================================================= +Standard_Boolean OpenGl_ShaderProgram::SetAttribute (const Handle(OpenGl_Context)& theCtx, + const GLchar* theName, + const OpenGl_Vec4& theValue) +{ + return SetAttribute (theCtx, GetAttributeLocation (theCtx, theName), theValue); +} + +// ======================================================================= +// function : SetAttribute +// purpose : +// ======================================================================= +Standard_Boolean OpenGl_ShaderProgram::SetAttribute (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + const OpenGl_Vec4& theValue) +{ + if (myProgramID == NO_PROGRAM || theIndex == INVALID_LOCATION) + { + return Standard_False; + } + + theCtx->core20fwd->glVertexAttrib4fv (theIndex, theValue); + return Standard_True; +} + // ======================================================================= // function : SetUniform // purpose : Specifies the value of the integer uniform variable @@ -1131,8 +1257,11 @@ void OpenGl_ShaderProgram::Release (const OpenGl_Context* theCtx) for (OpenGl_ShaderList::Iterator anIter (myShaderObjects); anIter.More(); anIter.Next()) { - anIter.ChangeValue()->Release (theCtx); - anIter.ChangeValue().Nullify(); + if (!anIter.Value().IsNull()) + { + anIter.ChangeValue()->Release (theCtx); + anIter.ChangeValue().Nullify(); + } } if (theCtx->core20 != NULL diff --git a/src/OpenGl/OpenGl_ShaderProgram.hxx b/src/OpenGl/OpenGl_ShaderProgram.hxx index bfa68c1fd1..959ad5d268 100755 --- a/src/OpenGl/OpenGl_ShaderProgram.hxx +++ b/src/OpenGl/OpenGl_ShaderProgram.hxx @@ -128,6 +128,7 @@ const int MaxStateTypes = 6; //! Wrapper for OpenGL program object. class OpenGl_ShaderProgram : public OpenGl_Resource { + friend class OpenGl_Workspace; public: @@ -260,6 +261,53 @@ public: GLint theIndex, OpenGl_Vec4& theValue) const; +public: + + //! Wrapper for glBindAttribLocation() + Standard_EXPORT Standard_Boolean SetAttributeName (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + const GLchar* theName); + + //! Wrapper for glVertexAttrib1f() + Standard_EXPORT Standard_Boolean SetAttribute (const Handle(OpenGl_Context)& theCtx, + const GLchar* theName, + GLfloat theValue); + + //! Wrapper for glVertexAttrib1f() + Standard_EXPORT Standard_Boolean SetAttribute (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + GLfloat theValue); + + //! Wrapper for glVertexAttrib2fv() + Standard_EXPORT Standard_Boolean SetAttribute (const Handle(OpenGl_Context)& theCtx, + const GLchar* theName, + const OpenGl_Vec2& theValue); + + //! Wrapper for glVertexAttrib2fv() + Standard_EXPORT Standard_Boolean SetAttribute (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + const OpenGl_Vec2& theValue); + + //! Wrapper for glVertexAttrib3fv() + Standard_EXPORT Standard_Boolean SetAttribute (const Handle(OpenGl_Context)& theCtx, + const GLchar* theName, + const OpenGl_Vec3& theValue); + + //! Wrapper for glVertexAttrib3fv() + Standard_EXPORT Standard_Boolean SetAttribute (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + const OpenGl_Vec3& theValue); + + //! Wrapper for glVertexAttrib4fv() + Standard_EXPORT Standard_Boolean SetAttribute (const Handle(OpenGl_Context)& theCtx, + const GLchar* theName, + const OpenGl_Vec4& theValue); + + //! Wrapper for glVertexAttrib4fv() + Standard_EXPORT Standard_Boolean SetAttribute (const Handle(OpenGl_Context)& theCtx, + GLint theIndex, + const OpenGl_Vec4& theValue); + public: //! Specifies the value of the integer uniform variable. diff --git a/src/OpenGl/OpenGl_Structure.cxx b/src/OpenGl/OpenGl_Structure.cxx index f491c42544..b0d8200478 100644 --- a/src/OpenGl/OpenGl_Structure.cxx +++ b/src/OpenGl/OpenGl_Structure.cxx @@ -13,10 +13,6 @@ // Alternatively, this file may be used under the terms of Open CASCADE // commercial license or contractual agreement. -#ifdef HAVE_CONFIG_H - #include -#endif - #include #include #include @@ -145,13 +141,11 @@ OpenGl_Structure::OpenGl_Structure (const Handle(Graphic3d_StructureManager)& th myAspectText(NULL), myHighlightColor(NULL), myNamedStatus(0), - myZLayer(0) + myZLayer(0), + myIsRaytracable (Standard_False), + myModificationState (0) { UpdateNamedStatus(); -#if HAVE_OPENCL - myIsRaytracable = Standard_False; - myModificationState = 0; -#endif } // ======================================================================= @@ -199,12 +193,10 @@ void OpenGl_Structure::UpdateTransformation() matcpy (myTransformation->mat, &Graphic3d_CStructure::Transformation[0][0]); -#ifdef HAVE_OPENCL if (myIsRaytracable) { UpdateStateWithAncestorStructures(); } -#endif } // ======================================================================= @@ -247,12 +239,10 @@ void OpenGl_Structure::SetAspectFace (const CALL_DEF_CONTEXTFILLAREA& theAspect) } myAspectFace->SetAspect (theAspect); -#ifdef HAVE_OPENCL if (myIsRaytracable) { UpdateStateWithAncestorStructures(); } -#endif } // ======================================================================= @@ -382,16 +372,12 @@ void OpenGl_Structure::UpdateNamedStatus() if (highlight) myNamedStatus |= OPENGL_NS_HIGHLIGHT; if (!visible) myNamedStatus |= OPENGL_NS_HIDE; -#ifdef HAVE_OPENCL if (myIsRaytracable) { UpdateStateWithAncestorStructures(); } -#endif } -#ifdef HAVE_OPENCL - // ======================================================================= // function : RegisterAncestorStructure // purpose : @@ -494,8 +480,6 @@ void OpenGl_Structure::SetRaytracableWithAncestorStructures() const } } -#endif - // ======================================================================= // function : Connect // purpose : @@ -506,7 +490,6 @@ void OpenGl_Structure::Connect (Graphic3d_CStructure& theStructure) Disconnect (theStructure); myConnected.Append (aStruct); -#ifdef HAVE_OPENCL if (aStruct->IsRaytracable()) { UpdateStateWithAncestorStructures(); @@ -514,7 +497,6 @@ void OpenGl_Structure::Connect (Graphic3d_CStructure& theStructure) } aStruct->RegisterAncestorStructure (this); -#endif } // ======================================================================= @@ -531,7 +513,6 @@ void OpenGl_Structure::Disconnect (Graphic3d_CStructure& theStructure) { myConnected.Remove (anIter); -#ifdef HAVE_OPENCL if (aStruct->IsRaytracable()) { UpdateStateWithAncestorStructures(); @@ -539,7 +520,6 @@ void OpenGl_Structure::Disconnect (Graphic3d_CStructure& theStructure) } aStruct->UnregisterAncestorStructure (this); -#endif return; } } @@ -574,13 +554,11 @@ void OpenGl_Structure::RemoveGroup (const Handle(Graphic3d_Group)& theGroup) { theGroup->Clear (Standard_False); - #ifdef HAVE_OPENCL if (((OpenGl_Group* )theGroup.operator->())->IsRaytracable()) { UpdateStateWithAncestorStructures(); UpdateRaytracableWithAncestorStructures(); } - #endif myGroups.Remove (aGroupIter); return; @@ -603,29 +581,23 @@ void OpenGl_Structure::Clear() // ======================================================================= void OpenGl_Structure::Clear (const Handle(OpenGl_Context)& theGlCtx) { -#ifdef HAVE_OPENCL Standard_Boolean aRaytracableGroupDeleted (Standard_False); -#endif // Release groups for (OpenGl_Structure::GroupIterator aGroupIter (myGroups); aGroupIter.More(); aGroupIter.Next()) { - #ifdef HAVE_OPENCL aRaytracableGroupDeleted |= aGroupIter.Value()->IsRaytracable(); - #endif // Delete objects aGroupIter.ChangeValue()->Release (theGlCtx); } myGroups.Clear(); -#ifdef HAVE_OPENCL if (aRaytracableGroupDeleted) { UpdateStateWithAncestorStructures(); UpdateRaytracableWithAncestorStructures(); } -#endif } // ======================================================================= @@ -846,10 +818,8 @@ void OpenGl_Structure::Release (const Handle(OpenGl_Context)& theGlCtx) OpenGl_Element::Destroy (theGlCtx, myAspectText); clearHighlightColor (theGlCtx); -#ifdef HAVE_OPENCL // Remove from connected list of ancestor UnregisterFromAncestorStructure(); -#endif } // ======================================================================= diff --git a/src/OpenGl/OpenGl_Structure.hxx b/src/OpenGl/OpenGl_Structure.hxx index 6b142a9369..fd7b3673ca 100644 --- a/src/OpenGl/OpenGl_Structure.hxx +++ b/src/OpenGl/OpenGl_Structure.hxx @@ -160,8 +160,6 @@ public: //! Returns OpenGL persistent translation. const TEL_TRANSFORM_PERSISTENCE* PersistentTranslation() const { return myTransPers; } -#ifdef HAVE_OPENCL - //! Returns structure modification state (for ray-tracing). Standard_Size ModificationState() const { return myModificationState; } @@ -171,14 +169,10 @@ public: //! Is the structure ray-tracable (contains ray-tracable elements)? Standard_Boolean IsRaytracable() const { return myIsRaytracable; } -#endif - protected: Standard_EXPORT virtual ~OpenGl_Structure(); -#ifdef HAVE_OPENCL - //! Registers ancestor connected structure (for updating ray-tracing state). void RegisterAncestorStructure (const OpenGl_Structure* theStructure) const; @@ -197,8 +191,6 @@ protected: //! Sets ray-tracable status for structure and its parents. void SetRaytracableWithAncestorStructures() const; -#endif - protected: OpenGl_Matrix* myTransformation; @@ -216,11 +208,9 @@ protected: OpenGl_ListOfStructure myConnected; -#ifdef HAVE_OPENCL mutable OpenGl_ListOfStructure myAncestorStructures; mutable Standard_Boolean myIsRaytracable; mutable Standard_Size myModificationState; -#endif public: diff --git a/src/OpenGl/OpenGl_TextureBufferArb.cxx b/src/OpenGl/OpenGl_TextureBufferArb.cxx index 4f068285d6..879bb34ebb 100644 --- a/src/OpenGl/OpenGl_TextureBufferArb.cxx +++ b/src/OpenGl/OpenGl_TextureBufferArb.cxx @@ -127,6 +127,45 @@ bool OpenGl_TextureBufferArb::Init (const Handle(OpenGl_Context)& theGlCtx, return true; } +// ======================================================================= +// function : Init +// purpose : +// ======================================================================= +bool OpenGl_TextureBufferArb::Init (const Handle(OpenGl_Context)& theGlCtx, + const GLuint theComponentsNb, + const GLsizei theElemsNb, + const GLuint* theData) +{ + if (theComponentsNb != 1 + && theComponentsNb != 2 + && theComponentsNb != 3 + && theComponentsNb != 4) + { + // unsupported format + return false; + } + else if (!Create (theGlCtx) + || !OpenGl_VertexBuffer::Init (theGlCtx, theComponentsNb, theElemsNb, theData)) + { + return false; + } + + switch (theComponentsNb) + { + case 1: myTexFormat = GL_R32I; break; + case 2: myTexFormat = GL_RG32I; break; + case 3: myTexFormat = GL_RGB32I; break; + case 4: myTexFormat = GL_RGBA32I; break; + } + + Bind (theGlCtx); + BindTexture (theGlCtx); + theGlCtx->arbTBO->glTexBuffer (GetTarget(), myTexFormat, myBufferId); + UnbindTexture (theGlCtx); + Unbind (theGlCtx); + return true; +} + // ======================================================================= // function : BindTexture // purpose : diff --git a/src/OpenGl/OpenGl_TextureBufferArb.hxx b/src/OpenGl/OpenGl_TextureBufferArb.hxx index b6c08ddc30..42abd344b3 100644 --- a/src/OpenGl/OpenGl_TextureBufferArb.hxx +++ b/src/OpenGl/OpenGl_TextureBufferArb.hxx @@ -69,6 +69,13 @@ public: const GLsizei theElemsNb, const GLfloat* theData); + //! Perform TBO initialization with specified data. + //! Existing data will be deleted. + Standard_EXPORT bool Init (const Handle(OpenGl_Context)& theGlCtx, + const GLuint theComponentsNb, + const GLsizei theElemsNb, + const GLuint* theData); + //! Bind TBO to specified Texture Unit. Standard_EXPORT void BindTexture (const Handle(OpenGl_Context)& theGlCtx, const GLenum theTextureUnit = GL_TEXTURE0) const; diff --git a/src/OpenGl/OpenGl_View.cxx b/src/OpenGl/OpenGl_View.cxx index 6d4db6cb21..d3c006bf33 100644 --- a/src/OpenGl/OpenGl_View.cxx +++ b/src/OpenGl/OpenGl_View.cxx @@ -13,10 +13,6 @@ // Alternatively, this file may be used under the terms of Open CASCADE // commercial license or contractual agreement. -#ifdef HAVE_CONFIG_H - #include -#endif - #include #include @@ -97,9 +93,7 @@ OpenGl_View::OpenGl_View (const CALL_DEF_VIEWCONTEXT &AContext, myCurrLightSourceState = myStateCounter->Increment(); -#ifdef HAVE_OPENCL myModificationState = 1; // initial state -#endif } /*----------------------------------------------------------------------*/ @@ -145,18 +139,14 @@ void OpenGl_View::SetTextureEnv (const Handle(OpenGl_Context)& theCtx, if (!anImage.IsNull()) myTextureEnv->Init (theCtx, *anImage.operator->(), theTexture->Type()); -#ifdef HAVE_OPENCL myModificationState++; -#endif } void OpenGl_View::SetSurfaceDetail (const Visual3d_TypeOfSurfaceDetail theMode) { mySurfaceDetail = theMode; -#ifdef HAVE_OPENCL myModificationState++; -#endif } // ======================================================================= diff --git a/src/OpenGl/OpenGl_View.hxx b/src/OpenGl/OpenGl_View.hxx index a88fdaf4c0..ef3dd2c7eb 100644 --- a/src/OpenGl/OpenGl_View.hxx +++ b/src/OpenGl/OpenGl_View.hxx @@ -171,7 +171,7 @@ class OpenGl_View : public MMgt_TShared const Aspect_CLayer2d& theCOverLayer); - void DrawBackground (const Handle(OpenGl_Workspace)& theWorkspace); + void DrawBackground (OpenGl_Workspace& theWorkspace); //! Returns list of OpenGL Z-layers. const OpenGl_LayerList& LayerList() const { return myZLayers; } @@ -194,10 +194,8 @@ class OpenGl_View : public MMgt_TShared return myImmediateList; } -#ifdef HAVE_OPENCL //! Returns modification state for ray-tracing. Standard_Size ModificationState() const { return myModificationState; } -#endif protected: @@ -269,9 +267,7 @@ protected: StateInfo myLastViewMappingState; StateInfo myLastLightSourceState; -#ifdef HAVE_OPENCL Standard_Size myModificationState; -#endif public: diff --git a/src/OpenGl/OpenGl_View_2.cxx b/src/OpenGl/OpenGl_View_2.cxx index 46014e1d1a..a53e7a433c 100644 --- a/src/OpenGl/OpenGl_View_2.cxx +++ b/src/OpenGl/OpenGl_View_2.cxx @@ -157,13 +157,13 @@ static void bind_light (const OpenGl_Light& theLight, /*----------------------------------------------------------------------*/ -void OpenGl_View::DrawBackground (const Handle(OpenGl_Workspace)& theWorkspace) +void OpenGl_View::DrawBackground (OpenGl_Workspace& theWorkspace) { - if ( (theWorkspace->NamedStatus & OPENGL_NS_WHITEBACK) == 0 && + if ( (theWorkspace.NamedStatus & OPENGL_NS_WHITEBACK) == 0 && ( myBgTexture.TexId != 0 || myBgGradient.type != Aspect_GFM_NONE ) ) { - const Standard_Integer aViewWidth = theWorkspace->Width(); - const Standard_Integer aViewHeight = theWorkspace->Height(); + const Standard_Integer aViewWidth = theWorkspace.Width(); + const Standard_Integer aViewHeight = theWorkspace.Height(); glPushAttrib( GL_ENABLE_BIT | GL_TEXTURE_BIT ); @@ -319,7 +319,7 @@ void OpenGl_View::DrawBackground (const Handle(OpenGl_Workspace)& theWorkspace) glDisable( GL_BLEND ); //push GL_ENABLE_BIT - glColor3fv (theWorkspace->BackgroundColor().rgb); + glColor3fv (theWorkspace.BackgroundColor().rgb); glTexEnvi (GL_TEXTURE_ENV, GL_TEXTURE_ENV_MODE, GL_DECAL); //push GL_TEXTURE_BIT // Note that texture is mapped using GL_REPEAT wrapping mode so integer part @@ -340,11 +340,11 @@ void OpenGl_View::DrawBackground (const Handle(OpenGl_Workspace)& theWorkspace) glPopAttrib(); //GL_ENABLE_BIT | GL_TEXTURE_BIT - if (theWorkspace->UseZBuffer()) + if (theWorkspace.UseZBuffer()) glEnable (GL_DEPTH_TEST); /* GL_DITHER on/off pour le trace */ - if (theWorkspace->Dither()) + if (theWorkspace.Dither()) glEnable (GL_DITHER); else glDisable (GL_DITHER); @@ -439,7 +439,7 @@ void OpenGl_View::Render (const Handle(OpenGl_PrinterContext)& thePrintContext, // ==================================== // Render background - DrawBackground (theWorkspace); + DrawBackground (*theWorkspace); // Switch off lighting by default glDisable(GL_LIGHTING); diff --git a/src/OpenGl/OpenGl_Workspace.cxx b/src/OpenGl/OpenGl_Workspace.cxx index 8f08b4b0d6..540c1e261f 100644 --- a/src/OpenGl/OpenGl_Workspace.cxx +++ b/src/OpenGl/OpenGl_Workspace.cxx @@ -146,6 +146,12 @@ OpenGl_Workspace::OpenGl_Workspace (const Handle(OpenGl_Display)& theDisplay, NamedStatus (0), HighlightColor (&THE_WHITE_COLOR), // + myComputeInitStatus (OpenGl_RT_NONE), + myIsRaytraceDataValid (Standard_False), + myTraversalStackSize (THE_DEFAULT_STACK_SIZE), + myViewModificationStatus (0), + myLayersModificationStatus (0), + // myTransientDrawToFront (Standard_True), myBackBufferRestored (Standard_False), myIsImmediateDrawn (Standard_False), @@ -186,18 +192,6 @@ OpenGl_Workspace::OpenGl_Workspace (const Handle(OpenGl_Display)& theDisplay, // Polygon Offset EnablePolygonOffset(); - -#ifdef HAVE_OPENCL - - myComputeInitStatus = OpenGl_CLIS_NONE; - - myViewModificationStatus = 0; - myLayersModificationStatus = 0; - - myIsRaytraceDataValid = Standard_False; - myToUpdateRaytraceData = Standard_False; - -#endif } // ======================================================================= @@ -217,9 +211,7 @@ Standard_Boolean OpenGl_Workspace::SetImmediateModeDrawToFront (const Standard_B // ======================================================================= OpenGl_Workspace::~OpenGl_Workspace() { -#ifdef HAVE_OPENCL - ReleaseOpenCL(); -#endif + ReleaseRaytraceResources(); } // ======================================================================= @@ -567,10 +559,8 @@ void OpenGl_Workspace::Redraw (const Graphic3d_CView& theCView, toSwap = 0; // no need to swap buffers } -#ifdef HAVE_OPENCL - if (!theCView.IsRaytracing || myComputeInitStatus == OpenGl_CLIS_FAIL) + if (!theCView.IsRaytracing || myComputeInitStatus == OpenGl_RT_FAIL) { -#endif const Standard_Boolean isImmediate = !myView->ImmediateStructures().IsEmpty(); redraw1 (theCView, theCUnderLayer, theCOverLayer, isImmediate ? 0 : toSwap); if (isImmediate) @@ -579,18 +569,16 @@ void OpenGl_Workspace::Redraw (const Graphic3d_CView& theCView, } theCView.WasRedrawnGL = Standard_True; -#ifdef HAVE_OPENCL } else { int aSizeX = aFrameBuffer != NULL ? aFrameBuffer->GetVPSizeX() : myWidth; int aSizeY = aFrameBuffer != NULL ? aFrameBuffer->GetVPSizeY() : myHeight; - Raytrace (theCView, aSizeX, aSizeY, toSwap); + Raytrace (theCView, aSizeX, aSizeY, toSwap, aFrameBuffer); theCView.WasRedrawnGL = Standard_False; } -#endif if (aFrameBuffer != NULL) { diff --git a/src/OpenGl/OpenGl_Workspace.hxx b/src/OpenGl/OpenGl_Workspace.hxx index dac1880ea3..d2dbc07804 100755 --- a/src/OpenGl/OpenGl_Workspace.hxx +++ b/src/OpenGl/OpenGl_Workspace.hxx @@ -16,12 +16,8 @@ #ifndef _OpenGl_Workspace_Header #define _OpenGl_Workspace_Header -#ifdef HAVE_OPENCL - #include - #include - - #include -#endif +#include +#include #include #include @@ -46,12 +42,11 @@ #include #include +#include #include #include #include -#ifdef HAVE_OPENCL - #include -#endif +#include #include #include #include @@ -59,6 +54,10 @@ #include #include +#include +#include +#include + class OpenGl_AspectLine; class OpenGl_AspectMarker; class OpenGl_AspectText; @@ -233,24 +232,119 @@ protected: void setTextureParams (Handle(OpenGl_Texture)& theTexture, const Handle(Graphic3d_TextureParams)& theParams); -#ifdef HAVE_OPENCL - -public: - - //! Returns information about OpenCL device used for computations. - Standard_Boolean GetOpenClDeviceInfo ( - NCollection_DataMap& theInfo) const; - protected: - //! Describes result of OpenCL initializing. - enum OpenClInitStatus + //! Result of OpenGL shaders initialization. + enum RaytraceInitStatus { - OpenGl_CLIS_NONE, - OpenGl_CLIS_INIT, - OpenGl_CLIS_FAIL + OpenGl_RT_NONE, + OpenGl_RT_INIT, + OpenGl_RT_FAIL }; + //! Defines frequently used shader variables. + enum ShaderVariableIndex + { + OpenGl_RT_aPosition, + + OpenGl_RT_uOriginLT, + OpenGl_RT_uOriginLB, + OpenGl_RT_uOriginRT, + OpenGl_RT_uOriginRB, + + OpenGl_RT_uDirectLT, + OpenGl_RT_uDirectLB, + OpenGl_RT_uDirectRT, + OpenGl_RT_uDirectRB, + + OpenGl_RT_uSceneRad, + OpenGl_RT_uSceneEps, + + OpenGl_RT_uLightAmbnt, + OpenGl_RT_uLightCount, + + OpenGl_RT_uShadEnabled, + OpenGl_RT_uReflEnabled, + + OpenGl_RT_uInputTexture, + + OpenGl_RT_uOffsetX, + OpenGl_RT_uOffsetY, + OpenGl_RT_uSamples, + + OpenGl_RT_NbVariables // special field + }; + + //! Defines texture samplers. + enum ShaderSamplerNames + { + OpenGl_RT_SceneNodeInfoTexture = 0, + OpenGl_RT_SceneMinPointTexture = 1, + OpenGl_RT_SceneMaxPointTexture = 2, + + OpenGl_RT_ObjectNodeInfoTexture = 3, + OpenGl_RT_ObjectMinPointTexture = 4, + OpenGl_RT_ObjectMaxPointTexture = 5, + + OpenGl_RT_GeometryVertexTexture = 6, + OpenGl_RT_GeometryNormalTexture = 7, + OpenGl_RT_GeometryTriangTexture = 8, + + OpenGl_RT_EnvironmentMapTexture = 9, + + OpenGl_RT_RaytraceMaterialTexture = 10, + OpenGl_RT_RaytraceLightSrcTexture = 11, + + OpenGl_RT_FSAAInputTexture = 12 + }; + + //! Tool class for management of shader sources. + class ShaderSource + { + public: + + //! Creates new uninitialized shader source. + ShaderSource() + { + // + } + + //! Creates new shader source from specified file. + ShaderSource (const TCollection_AsciiString& theFileName) + { + Load (&theFileName, 1); + } + + public: + + //! Returns prefix to insert before the source. + const TCollection_AsciiString& Prefix() const + { + return myPrefix; + } + + //! Sets prefix to insert before the source. + void SetPrefix (const TCollection_AsciiString& thePrefix) + { + myPrefix = thePrefix; + } + + //! Returns shader source combined with prefix. + TCollection_AsciiString Source() const; + + //! Loads shader source from specified files. + void Load (const TCollection_AsciiString* theFileNames, const Standard_Integer theCount); + + private: + + TCollection_AsciiString mySource; //!< Source string of the shader object + TCollection_AsciiString myPrefix; //!< Prefix to insert before the source + + }; + + //! Default size of traversal stack. + static const Standard_Integer THE_DEFAULT_STACK_SIZE = 24; + protected: //! @name methods related to ray-tracing //! Updates 3D scene geometry for ray-tracing. @@ -301,47 +395,53 @@ protected: //! @name methods related to ray-tracing Standard_Boolean AddRaytracePolygonArray (OpenGl_TriangleSet* theSet, const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID); - //! Initializes OpenCL resources. - Standard_Boolean InitOpenCL(); + //! Loads and compiles shader object from specified source. + Handle(OpenGl_ShaderObject) LoadShader (const ShaderSource& theSource, GLenum theType); - //! Releases OpenCL resources. - void ReleaseOpenCL(); + //! Performs safe exit when shaders initialization fails. + Standard_Boolean SafeFailBack (const TCollection_ExtendedString& theMessage); - //! Resizes OpenCL output image. - Standard_Boolean ResizeRaytraceOutputBuffer (const cl_int theSizeX, const cl_int theSizeY); + //! Initializes OpenGL/GLSL shader programs. + Standard_Boolean InitRaytraceResources(); - //! Writes scene geometry to OpenCl device. - Standard_Boolean WriteRaytraceSceneToDevice(); + //! Releases OpenGL/GLSL shader programs. + void ReleaseRaytraceResources(); - //! Runs OpenCL ray-tracing kernels. - Standard_Boolean RunRaytraceOpenCLKernelsOld (const Graphic3d_CView& theCView, - const GLfloat theOrigins[16], - const GLfloat theDirects[16], - const int theSizeX, - const int theSizeY); + //! Uploads ray-trace data to the GPU. + Standard_Boolean UploadRaytraceData(); - //! Launches OpenCL ray-tracing kernels. - Standard_Boolean RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView, - const Standard_ShortReal theOrigins[16], - const Standard_ShortReal theDirects[16], - const Standard_Integer theSizeX, - const Standard_Integer theSizeY); + //! Resizes OpenGL frame buffers. + Standard_Boolean ResizeRaytraceBuffers (const Standard_Integer theSizeX, + const Standard_Integer theSizeY); - //! Redraws the window using OpenCL ray tracing. + //! Generates viewing rays for corners of screen quad. + void UpdateCamera (const NCollection_Mat4& theOrientation, + const NCollection_Mat4& theViewMapping, + OpenGl_Vec3 theOrigins[4], + OpenGl_Vec3 theDirects[4]); + + //! Runs ray-tracing shader programs. + Standard_Boolean RunRaytraceShaders (const Graphic3d_CView& theCView, + const Standard_Integer theSizeX, + const Standard_Integer theSizeY, + const OpenGl_Vec3 theOrigins[4], + const OpenGl_Vec3 theDirects[4], + OpenGl_FrameBuffer* theFrameBuffer); + + //! Redraws the window using OpenGL/GLSL ray-tracing. Standard_Boolean Raytrace (const Graphic3d_CView& theCView, - const int theSizeX, const int theSizeY, const Tint theToSwap); + const Standard_Integer theSizeX, + const Standard_Integer theSizeY, + const Standard_Boolean theToSwap, + OpenGl_FrameBuffer* theFrameBuffer); protected: //! @name fields related to ray-tracing - //! Result of OpenCL initialization. - OpenClInitStatus myComputeInitStatus; - //! Is ATI/AMD OpenCL platform used? - Standard_Boolean myIsAmdComputePlatform; + //! Result of shaders initialization. + RaytraceInitStatus myComputeInitStatus; //! Is geometry data valid? Standard_Boolean myIsRaytraceDataValid; - //! Is geometry data musty be updated? - Standard_Boolean myToUpdateRaytraceData; //! 3D scene geometry data for ray-tracing. OpenGl_RaytraceGeometry myRaytraceGeometry; @@ -351,54 +451,57 @@ protected: //! @name fields related to ray-tracing //! Scene epsilon to prevent self-intersections. Standard_ShortReal myRaytraceSceneEpsilon; - //! OpenCL context. - cl_context myComputeContext; - //! OpenCL command queue. - cl_command_queue myComputeQueue; - //! OpenCL computing program. - cl_program myRaytraceProgram; - //! OpenCL ray-tracing render kernel. - cl_kernel myRaytraceRenderKernel; - //! OpenCL adaptive anti-aliasing kernel. - cl_kernel myRaytraceSmoothKernel; + //! Actual size of traversal stack in shader program. + Standard_Integer myTraversalStackSize; - //! OpenCL image to store environment map. - cl_mem myRaytraceEnvironment; - //! OpenCL image to store rendering result. - cl_mem myRaytraceOutputImage; - //! OpenCL image to store anti-aliasing result. - cl_mem myRaytraceOutputImageAA; + //! OpenGL/GLSL source of ray-tracing fragment shader. + ShaderSource myRaytraceShaderSource; + //! OpenGL/GLSL source of adaptive-AA fragment shader. + ShaderSource myPostFSAAShaderSource; - //! OpenGL texture to store rendering result. - Handle(OpenGl_Texture) myRaytraceOutputTexture; - //! OpenGL texture to store anti-aliasing result. - Handle(OpenGl_Texture) myRaytraceOutputTextureAA; + //! OpenGL/GLSL ray-tracing fragment shader. + Handle(OpenGl_ShaderObject) myRaytraceShader; + //! OpenGL/GLSL adaptive-AA fragment shader. + Handle(OpenGl_ShaderObject) myPostFSAAShader; - //! OpenCL buffer of material properties. - cl_mem myRaytraceMaterialBuffer; - //! OpenCL buffer of light source properties. - cl_mem myRaytraceLightSourceBuffer; + //! OpenGL/GLSL ray-tracing shader program. + Handle(OpenGl_ShaderProgram) myRaytraceProgram; + //! OpenGL/GLSL adaptive-AA shader program. + Handle(OpenGl_ShaderProgram) myPostFSAAProgram; - //! OpenCL buffer of vertex coords. - cl_mem myGeometryVertexBuffer; - //! OpenCL buffer of vertex normals. - cl_mem myGeometryNormalBuffer; - //! OpenCL buffer of triangle indices. - cl_mem myGeometryTriangBuffer; + //! Texture buffer of data records of high-level BVH nodes. + Handle(OpenGl_TextureBufferArb) mySceneNodeInfoTexture; + //! Texture buffer of minimum points of high-level BVH nodes. + Handle(OpenGl_TextureBufferArb) mySceneMinPointTexture; + //! Texture buffer of maximum points of high-level BVH nodes. + Handle(OpenGl_TextureBufferArb) mySceneMaxPointTexture; - //! OpenCL buffer of data records of high-level BVH nodes. - cl_mem mySceneNodeInfoBuffer; - //! OpenCL buffer of minimum points of high-level BVH nodes. - cl_mem mySceneMinPointBuffer; - //! OpenCL buffer of maximum points of high-level BVH nodes. - cl_mem mySceneMaxPointBuffer; + //! Texture buffer of data records of bottom-level BVH nodes. + Handle(OpenGl_TextureBufferArb) myObjectNodeInfoTexture; + //! Texture buffer of minimum points of bottom-level BVH nodes. + Handle(OpenGl_TextureBufferArb) myObjectMinPointTexture; + //! Texture buffer of maximum points of bottom-level BVH nodes. + Handle(OpenGl_TextureBufferArb) myObjectMaxPointTexture; - //! OpenCL buffer of data records of bottom-level BVH nodes. - cl_mem myObjectNodeInfoBuffer; - //! OpenCL buffer of minimum points of bottom-level BVH nodes. - cl_mem myObjectMinPointBuffer; - //! OpenCL buffer of maximum points of bottom-level BVH nodes. - cl_mem myObjectMaxPointBuffer; + //! Texture buffer of vertex coords. + Handle(OpenGl_TextureBufferArb) myGeometryVertexTexture; + //! Texture buffer of vertex normals. + Handle(OpenGl_TextureBufferArb) myGeometryNormalTexture; + //! Texture buffer of triangle indices. + Handle(OpenGl_TextureBufferArb) myGeometryTriangTexture; + + //! Texture buffer of material properties. + Handle(OpenGl_TextureBufferArb) myRaytraceMaterialTexture; + //! Texture buffer of light source properties. + Handle(OpenGl_TextureBufferArb) myRaytraceLightSrcTexture; + + //! Vertex buffer (VBO) for drawing dummy quad. + OpenGl_VertexBuffer myRaytraceScreenQuad; + + //! Framebuffer (FBO) to perform adaptive FSAA. + Handle(OpenGl_FrameBuffer) myRaytraceFBO1; + //! Framebuffer (FBO) to perform adaptive FSAA. + Handle(OpenGl_FrameBuffer) myRaytraceFBO2; //! State of OpenGL view. Standard_Size myViewModificationStatus; @@ -408,7 +511,8 @@ protected: //! @name fields related to ray-tracing //! State of OpenGL structures reflected to ray-tracing. std::map myStructureStates; -#endif // HAVE_OPENCL + //! Cached locations of frequently used uniform variables. + Standard_Integer myUniformLocations[2][OpenGl_RT_NbVariables]; protected: //! @name protected fields diff --git a/src/OpenGl/OpenGl_Workspace_Raytrace.cxx b/src/OpenGl/OpenGl_Workspace_Raytrace.cxx index 2090cefcda..5074174740 100755 --- a/src/OpenGl/OpenGl_Workspace_Raytrace.cxx +++ b/src/OpenGl/OpenGl_Workspace_Raytrace.cxx @@ -13,46 +13,26 @@ // Alternatively, this file may be used under the terms of Open CASCADE // commercial license or contractual agreement. -#ifdef HAVE_CONFIG_H - #include -#endif - -#ifdef HAVE_OPENCL - -#include - -#if defined(_WIN32) - - #include - #include - - #pragma comment (lib, "DelayImp.lib") - #pragma comment (lib, "OpenCL.lib") - -#elif defined(__APPLE__) && !defined(MACOSX_USE_GLX) - #include -#else - #include -#endif - -#include +#include +#include +#include #include +#include #include #include +#include +#include #include using namespace OpenGl_Raytrace; //! Use this macro to output ray-tracing debug info -//#define RAY_TRACE_PRINT_INFO +// #define RAY_TRACE_PRINT_INFO #ifdef RAY_TRACE_PRINT_INFO #include #endif -//! OpenCL source of ray-tracing kernels. -extern const char THE_RAY_TRACE_OPENCL_SOURCE[]; - // ======================================================================= // function : MatVecMult // purpose : Multiples 4x4 matrix by 4D vector @@ -71,86 +51,6 @@ BVH_Vec4f MatVecMult (const T m[16], const BVH_Vec4f& v) m[11] * v.z() + m[15] * v.w())); } -// ======================================================================= -// function : UpdateRaytraceEnvironmentMap -// purpose : Updates environment map for ray-tracing -// ======================================================================= -Standard_Boolean OpenGl_Workspace::UpdateRaytraceEnvironmentMap() -{ - if (myView.IsNull()) - return Standard_False; - - if (myViewModificationStatus == myView->ModificationState()) - return Standard_True; - - cl_int anError = CL_SUCCESS; - - if (myRaytraceEnvironment != NULL) - clReleaseMemObject (myRaytraceEnvironment); - - Standard_Integer aSizeX = 1; - Standard_Integer aSizeY = 1; - - if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE) - { - aSizeX = (myView->TextureEnv()->SizeX() <= 0) ? 1 : myView->TextureEnv()->SizeX(); - aSizeY = (myView->TextureEnv()->SizeY() <= 0) ? 1 : myView->TextureEnv()->SizeY(); - } - - cl_image_format anImageFormat; - - anImageFormat.image_channel_order = CL_RGBA; - anImageFormat.image_channel_data_type = CL_FLOAT; - - myRaytraceEnvironment = clCreateImage2D (myComputeContext, - CL_MEM_READ_ONLY, &anImageFormat, aSizeX, aSizeY, 0, NULL, &anError); - - cl_float* aPixelData = new cl_float[aSizeX * aSizeY * 4]; - - // Note: texture format is not compatible with OpenCL image - // (it's not possible to create image directly from texture) - - if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE) - { - myView->TextureEnv()->Bind (GetGlContext()); - - glGetTexImage (GL_TEXTURE_2D, - 0, - GL_RGBA, - GL_FLOAT, - aPixelData); - - myView->TextureEnv()->Unbind (GetGlContext()); - } - else - { - for (Standard_Integer aPixel = 0; aPixel < aSizeX * aSizeY * 4; ++aPixel) - aPixelData[aPixel] = 0.f; - } - - size_t anImageOffset[] = { 0, - 0, - 0 }; - - size_t anImageRegion[] = { aSizeX, - aSizeY, - 1 }; - - anError |= clEnqueueWriteImage (myComputeQueue, myRaytraceEnvironment, - CL_TRUE, anImageOffset, anImageRegion, 0, 0, aPixelData, 0, NULL, NULL); - -#ifdef RAY_TRACE_PRINT_INFO - if (anError != CL_SUCCESS) - std::cout << "Error! Failed to write environment map image!" << std::endl; -#endif - - delete[] aPixelData; - - myViewModificationStatus = myView->ModificationState(); - - return (anError == CL_SUCCESS); -} - // ======================================================================= // function : UpdateRaytraceGeometry // purpose : Updates 3D scene geometry for ray tracing @@ -263,10 +163,12 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theC myRaytraceSceneRadius = 2.f /* scale factor */ * Max (aMinRadius, aMaxRadius); - myRaytraceSceneEpsilon = Max (1e-4f, - myRaytraceGeometry.Box().Size().Length() * 1e-4f); + const BVH_Vec4f aSize = myRaytraceGeometry.Box().Size(); - return WriteRaytraceSceneToDevice(); + myRaytraceSceneEpsilon = Max (1e-4f, 1e-4f * sqrtf ( + aSize.x() * aSize.x() + aSize.y() * aSize.y() + aSize.z() * aSize.z())); + + return UploadRaytraceData(); } delete [] aTransform; @@ -345,10 +247,10 @@ void CreateMaterial (const OPENGL_SURF_PROP& theProp, OpenGl_RaytraceMaterial& t const float aReflectionScale = 0.75f / aMaxRefl; - theMaterial.Reflection = BVH_Vec4f (theProp.speccol.rgb[0] * theProp.spec, - theProp.speccol.rgb[1] * theProp.spec, - theProp.speccol.rgb[2] * theProp.spec, - 0.f) * aReflectionScale; + theMaterial.Reflection = BVH_Vec4f (theProp.speccol.rgb[0] * theProp.spec * aReflectionScale, + theProp.speccol.rgb[1] * theProp.spec * aReflectionScale, + theProp.speccol.rgb[2] * theProp.spec * aReflectionScale, + 0.f); } // ======================================================================= @@ -857,7 +759,7 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble th { myRaytraceGeometry.Sources.clear(); - myRaytraceGeometry.GlobalAmbient = BVH_Vec4f (0.0f, 0.0f, 0.0f, 0.0f); + myRaytraceGeometry.Ambient = BVH_Vec4f (0.0f, 0.0f, 0.0f, 0.0f); for (OpenGl_ListOfLight::Iterator anItl (myView->LightList()); anItl.More(); anItl.Next()) { @@ -865,10 +767,10 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble th if (aLight.Type == Visual3d_TOLS_AMBIENT) { - myRaytraceGeometry.GlobalAmbient += BVH_Vec4f (aLight.Color.r(), - aLight.Color.g(), - aLight.Color.b(), - 0.0f); + myRaytraceGeometry.Ambient += BVH_Vec4f (aLight.Color.r(), + aLight.Color.g(), + aLight.Color.b(), + 0.0f); continue; } @@ -893,562 +795,712 @@ Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble th if (aLight.IsHeadlight) aPosition = MatVecMult (theInvModelView, aPosition); + myRaytraceGeometry.Sources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition)); } - cl_int anError = CL_SUCCESS; - - if (myRaytraceLightSourceBuffer != NULL) - clReleaseMemObject (myRaytraceLightSourceBuffer); - - Standard_Integer aLightBufferSize = myRaytraceGeometry.Sources.size() != 0 ? - static_cast (myRaytraceGeometry.Sources.size()) : 1; - - myRaytraceLightSourceBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, - aLightBufferSize * sizeof(OpenGl_RaytraceLight), NULL, &anError); + if (myRaytraceLightSrcTexture.IsNull()) // create light source buffer + { + myRaytraceLightSrcTexture = new OpenGl_TextureBufferArb; + if (!myRaytraceLightSrcTexture->Create (myGlContext)) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error: Failed to create light source buffer" << std::endl; +#endif + return Standard_False; + } + } + if (myRaytraceGeometry.Sources.size() != 0) { - const void* aDataPtr = myRaytraceGeometry.Sources.front().Packed(); + const GLfloat* aDataPtr = myRaytraceGeometry.Sources.front().Packed(); - anError |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0, - aLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr, 0, NULL, NULL); - } + bool aResult = myRaytraceLightSrcTexture->Init ( + myGlContext, 4, myRaytraceGeometry.Sources.size() * 2, aDataPtr); + if (!aResult) + { #ifdef RAY_TRACE_PRINT_INFO - if (anError != CL_SUCCESS) - { - std::cout << "Error! Failed to set light sources"; - - return Standard_False; - } + std::cout << "Error: Failed to upload light source buffer" << std::endl; #endif + return Standard_False; + } + } return Standard_True; } // ======================================================================= -// function : CheckOpenCL -// purpose : Checks OpenCL dynamic library availability +// function : UpdateRaytraceEnvironmentMap +// purpose : Updates environment map for ray-tracing // ======================================================================= -Standard_Boolean CheckOpenCL() +Standard_Boolean OpenGl_Workspace::UpdateRaytraceEnvironmentMap() { -#if defined ( _WIN32 ) - - __try - { - cl_uint aNbPlatforms; - clGetPlatformIDs (0, NULL, &aNbPlatforms); - } - __except (EXCEPTION_EXECUTE_HANDLER) - { + if (myView.IsNull()) return Standard_False; - } -#endif + if (myViewModificationStatus == myView->ModificationState()) + return Standard_True; - return Standard_True; -} - -// ======================================================================= -// function : InitOpenCL -// purpose : Initializes OpenCL objects -// ======================================================================= -Standard_Boolean OpenGl_Workspace::InitOpenCL() -{ - if (myComputeInitStatus != OpenGl_CLIS_NONE) + for (Standard_Integer anIdx = 0; anIdx < 2; ++anIdx) { - return myComputeInitStatus == OpenGl_CLIS_INIT; - } + const Handle(OpenGl_ShaderProgram)& aProgram = + anIdx == 0 ? myRaytraceProgram : myPostFSAAProgram; - if (!CheckOpenCL()) - { - myComputeInitStatus = OpenGl_CLIS_FAIL; // fail to load OpenCL library - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, - 0, - GL_DEBUG_SEVERITY_HIGH_ARB, - "Failed to load OpenCL dynamic library!"); - return Standard_False; - } - - // Obtain the list of platforms available - cl_uint aNbPlatforms = 0; - cl_int anError = clGetPlatformIDs (0, NULL, &aNbPlatforms); - cl_platform_id* aPlatforms = (cl_platform_id* )alloca (aNbPlatforms * sizeof(cl_platform_id)); - anError |= clGetPlatformIDs (aNbPlatforms, aPlatforms, NULL); - if (anError != CL_SUCCESS - || aNbPlatforms == 0) - { - myComputeInitStatus = OpenGl_CLIS_FAIL; - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, - 0, - GL_DEBUG_SEVERITY_HIGH_ARB, - "No any OpenCL platform installed!"); - return Standard_False; - } - - // Note: We try to find NVIDIA or AMD platforms with GPU devices! - cl_platform_id aPrefPlatform = NULL; - for (cl_uint aPlatIter = 0; aPlatIter < aNbPlatforms; ++aPlatIter) - { - char aName[256]; - anError = clGetPlatformInfo (aPlatforms[aPlatIter], CL_PLATFORM_NAME, - sizeof(aName), aName, NULL); - if (anError != CL_SUCCESS) + if (!aProgram.IsNull()) { - continue; - } + aProgram->Bind (myGlContext); - if (strncmp (aName, "NVIDIA", strlen ("NVIDIA")) == 0) - { - aPrefPlatform = aPlatforms[aPlatIter]; - - // Use optimizations for NVIDIA GPUs - myIsAmdComputePlatform = Standard_False; - } - else if (strncmp (aName, "AMD", strlen ("AMD")) == 0) - { - aPrefPlatform = (aPrefPlatform == NULL) - ? aPlatforms[aPlatIter] - : aPrefPlatform; - - // Use optimizations for ATI/AMD platform - myIsAmdComputePlatform = Standard_True; - } - } - - if (aPrefPlatform == NULL) - { - aPrefPlatform = aPlatforms[0]; - } - - // Obtain the list of devices available in the selected platform - cl_uint aNbDevices = 0; - anError = clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU, - 0, NULL, &aNbDevices); - - cl_device_id* aDevices = (cl_device_id* )alloca (aNbDevices * sizeof(cl_device_id)); - anError |= clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU, - aNbDevices, aDevices, NULL); - if (anError != CL_SUCCESS) - { - myComputeInitStatus = OpenGl_CLIS_FAIL; - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, - 0, - GL_DEBUG_SEVERITY_HIGH_ARB, - "Failed to get OpenCL GPU device!"); - return Standard_False; - } - - // Note: Simply get first available GPU - cl_device_id aDevice = aDevices[0]; - - // detect old contexts - char aVerClStr[256]; - clGetDeviceInfo (aDevice, CL_DEVICE_VERSION, - sizeof(aVerClStr), aVerClStr, NULL); - aVerClStr[strlen ("OpenCL 1.0")] = '\0'; - const bool isVer10 = strncmp (aVerClStr, "OpenCL 1.0", strlen ("OpenCL 1.0")) == 0; - - // Create OpenCL context - cl_context_properties aCtxProp[] = - { - #if defined(__APPLE__) && !defined(MACOSX_USE_GLX) - CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, - (cl_context_properties )CGLGetShareGroup (CGLGetCurrentContext()), - #elif defined(_WIN32) - CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform, - CL_GL_CONTEXT_KHR, (cl_context_properties )wglGetCurrentContext(), - CL_WGL_HDC_KHR, (cl_context_properties )wglGetCurrentDC(), - #else - CL_GL_CONTEXT_KHR, (cl_context_properties )glXGetCurrentContext(), - CL_GLX_DISPLAY_KHR, (cl_context_properties )glXGetCurrentDisplay(), - CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform, - #endif - 0 - }; - - myComputeContext = clCreateContext (aCtxProp, - #if defined(__APPLE__) && !defined(MACOSX_USE_GLX) - 0, NULL, // device will be taken from GL context - #else - 1, &aDevice, - #endif - NULL, NULL, &anError); - if (anError != CL_SUCCESS) - { - myComputeInitStatus = OpenGl_CLIS_FAIL; - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, - 0, - GL_DEBUG_SEVERITY_HIGH_ARB, - "Failed to initialize OpenCL context!"); - return Standard_False; - } - - // Create OpenCL program - const char* aSources[] = - { - isVer10 ? "#define M_PI_F ( float )( 3.14159265359f )\n" : "", - THE_RAY_TRACE_OPENCL_SOURCE - }; - myRaytraceProgram = clCreateProgramWithSource (myComputeContext, 2, - aSources, NULL, &anError); - if (anError != CL_SUCCESS) - { - myComputeInitStatus = OpenGl_CLIS_FAIL; - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, - 0, - GL_DEBUG_SEVERITY_HIGH_ARB, - "Failed to create OpenCL ray-tracing program!"); - return Standard_False; - } - - anError = clBuildProgram (myRaytraceProgram, 0, - NULL, NULL, NULL, NULL); - { - // Fetch build log - size_t aLogLen = 0; - cl_int aResult = clGetProgramBuildInfo (myRaytraceProgram, aDevice, - CL_PROGRAM_BUILD_LOG, 0, NULL, &aLogLen); - - char* aBuildLog = (char* )alloca (aLogLen); - aResult |= clGetProgramBuildInfo (myRaytraceProgram, aDevice, - CL_PROGRAM_BUILD_LOG, aLogLen, aBuildLog, NULL); - if (aResult == CL_SUCCESS) - { - if (anError != CL_SUCCESS) + if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE) { - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, - 0, - GL_DEBUG_SEVERITY_HIGH_ARB, - aBuildLog); + myView->TextureEnv()->Bind ( + myGlContext, GL_TEXTURE0 + OpenGl_RT_EnvironmentMapTexture); + + aProgram->SetUniform (myGlContext, "uEnvironmentEnable", 1); } else { - #ifdef RAY_TRACE_PRINT_INFO - std::cout << aBuildLog << std::endl; - #endif + aProgram->SetUniform (myGlContext, "uEnvironmentEnable", 0); + } + + aProgram->SetSampler (myGlContext, + "uEnvironmentMapTexture", OpenGl_RT_EnvironmentMapTexture); + } + } + + OpenGl_ShaderProgram::Unbind (myGlContext); + + myViewModificationStatus = myView->ModificationState(); + + return Standard_True; +} + +// ======================================================================= +// function : Source +// purpose : Returns shader source combined with prefix +// ======================================================================= +TCollection_AsciiString OpenGl_Workspace::ShaderSource::Source() const +{ + static const TCollection_AsciiString aVersion = "#version 140"; + + if (myPrefix.IsEmpty()) + { + return aVersion + "\n" + mySource; + } + + return aVersion + "\n" + myPrefix + "\n" + mySource; +} + +// ======================================================================= +// function : Load +// purpose : Loads shader source from specified files +// ======================================================================= +void OpenGl_Workspace::ShaderSource::Load ( + const TCollection_AsciiString* theFileNames, const Standard_Integer theCount) +{ + mySource.Clear(); + + for (Standard_Integer anIndex = 0; anIndex < theCount; ++anIndex) + { + OSD_File aFile (theFileNames[anIndex]); + + Standard_ASSERT_RETURN (aFile.Exists(), + "Error: Failed to find shader source file", /* none */); + + aFile.Open (OSD_ReadOnly, OSD_Protection()); + + TCollection_AsciiString aSource; + + Standard_ASSERT_RETURN (aFile.IsOpen(), + "Error: Failed to open shader source file", /* none */); + + aFile.Read (aSource, (Standard_Integer) aFile.Size()); + + if (!aSource.IsEmpty()) + { + mySource += TCollection_AsciiString ("\n") + aSource; + } + + aFile.Close(); + } +} + +// ======================================================================= +// function : LoadShader +// purpose : Creates new shader object with specified source +// ======================================================================= +Handle(OpenGl_ShaderObject) OpenGl_Workspace::LoadShader (const ShaderSource& theSource, GLenum theType) +{ + Handle(OpenGl_ShaderObject) aShader = new OpenGl_ShaderObject (theType); + + if (!aShader->Create (myGlContext)) + { + const TCollection_ExtendedString aMessage = "Error: Failed to create shader object"; + + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); + + aShader->Release (myGlContext.operator->()); + + return Handle(OpenGl_ShaderObject)(); + } + + if (!aShader->LoadSource (myGlContext, theSource.Source())) + { + const TCollection_ExtendedString aMessage = "Error: Failed to set shader source"; + + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); + + aShader->Release (myGlContext.operator->()); + + return Handle(OpenGl_ShaderObject)(); + } + + TCollection_AsciiString aBuildLog; + + if (!aShader->Compile (myGlContext)) + { + if (aShader->FetchInfoLog (myGlContext, aBuildLog)) + { + const TCollection_ExtendedString aMessage = + TCollection_ExtendedString ("Error: Failed to compile shader object:\n") + aBuildLog; + +#ifdef RAY_TRACE_PRINT_INFO + std::cout << aBuildLog << std::endl; +#endif + + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); + } + + aShader->Release (myGlContext.operator->()); + + return Handle(OpenGl_ShaderObject)(); + } + +#ifdef RAY_TRACE_PRINT_INFO + if (aShader->FetchInfoLog (myGlContext, aBuildLog)) + { + if (!aBuildLog.IsEmpty()) + { + std::cout << aBuildLog << std::endl; + } + else + { + std::cout << "Info: shader build log is empty" << std::endl; + } + } +#endif + + return aShader; +} + +// ======================================================================= +// function : SafeFailBack +// purpose : Performs safe exit when shaders initialization fails +// ======================================================================= +Standard_Boolean OpenGl_Workspace::SafeFailBack (const TCollection_ExtendedString& theMessage) +{ + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, theMessage); + + myComputeInitStatus = OpenGl_RT_FAIL; + + ReleaseRaytraceResources(); + + return Standard_False; +} + +// ======================================================================= +// function : InitRaytraceResources +// purpose : Initializes OpenGL/GLSL shader programs +// ======================================================================= +Standard_Boolean OpenGl_Workspace::InitRaytraceResources() +{ + Standard_Boolean aToRebuildShaders = Standard_False; + + if (myComputeInitStatus == OpenGl_RT_INIT) + { + if (!myIsRaytraceDataValid) + return Standard_True; + + const Standard_Integer aRequiredStackSize = + myRaytraceGeometry.HighLevelTreeDepth() + myRaytraceGeometry.BottomLevelTreeDepth(); + + if (myTraversalStackSize < aRequiredStackSize) + { + myTraversalStackSize = Max (aRequiredStackSize, THE_DEFAULT_STACK_SIZE); + + aToRebuildShaders = Standard_True; + } + else + { + if (aRequiredStackSize < myTraversalStackSize) + { + if (myTraversalStackSize > THE_DEFAULT_STACK_SIZE) + { + myTraversalStackSize = Max (aRequiredStackSize, THE_DEFAULT_STACK_SIZE); + + aToRebuildShaders = Standard_True; + } + } + } + + if (aToRebuildShaders) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Info: Rebuild shaders with stack size: " << myTraversalStackSize << std::endl; +#endif + + TCollection_AsciiString aStackSizeStr = + TCollection_AsciiString ("#define STACK_SIZE ") + TCollection_AsciiString (myTraversalStackSize); + + myRaytraceShaderSource.SetPrefix (aStackSizeStr); + myPostFSAAShaderSource.SetPrefix (aStackSizeStr); + + if (!myRaytraceShader->LoadSource (myGlContext, myRaytraceShaderSource.Source()) + || !myPostFSAAShader->LoadSource (myGlContext, myPostFSAAShaderSource.Source())) + { + return Standard_False; + } + + if (!myRaytraceShader->Compile (myGlContext) + || !myPostFSAAShader->Compile (myGlContext)) + { + return Standard_False; + } + + if (!myRaytraceProgram->Link (myGlContext) + || !myPostFSAAProgram->Link (myGlContext)) + { + return Standard_False; } } } - if (anError != CL_SUCCESS) + if (myComputeInitStatus == OpenGl_RT_NONE) { - return Standard_False; - } - - // Create OpenCL ray tracing kernels - myRaytraceRenderKernel = clCreateKernel (myRaytraceProgram, "RaytraceRender", &anError); - if (anError != CL_SUCCESS) - { - myComputeInitStatus = OpenGl_CLIS_FAIL; - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, - 0, - GL_DEBUG_SEVERITY_HIGH_ARB, - "Failed to create OpenCL ray-tracing kernel!"); - return Standard_False; - } - - myRaytraceSmoothKernel = clCreateKernel (myRaytraceProgram, "RaytraceSmooth", &anError); - if (anError != CL_SUCCESS) - { - myComputeInitStatus = OpenGl_CLIS_FAIL; - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, - 0, - GL_DEBUG_SEVERITY_HIGH_ARB, - "Failed to create OpenCL ray-tracing kernel!"); - return Standard_False; - } - - // Create OpenCL command queue - // Note: For profiling set CL_QUEUE_PROFILING_ENABLE - cl_command_queue_properties aProps = CL_QUEUE_PROFILING_ENABLE; - - myComputeQueue = clCreateCommandQueue (myComputeContext, aDevice, aProps, &anError); - if (anError != CL_SUCCESS) - { - myComputeInitStatus = OpenGl_CLIS_FAIL; - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, - 0, - GL_DEBUG_SEVERITY_HIGH_ARB, - "Failed to create OpenCL command queue!"); - - return Standard_False; - } - - myComputeInitStatus = OpenGl_CLIS_INIT; // initialized in normal way - return Standard_True; -} - -// ======================================================================= -// function : GetOpenClDeviceInfo -// purpose : Returns information about device used for computations -// ======================================================================= -Standard_Boolean OpenGl_Workspace::GetOpenClDeviceInfo (NCollection_DataMap& theInfo) const -{ - theInfo.Clear(); - if (myComputeContext == NULL) - { - return Standard_False; - } - - size_t aDevicesSize = 0; - cl_int anError = clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, 0, NULL, &aDevicesSize); - cl_device_id* aDevices = (cl_device_id* )alloca (aDevicesSize); - anError |= clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, aDevicesSize, aDevices, NULL); - if (anError != CL_SUCCESS) - { - return Standard_False; - } - - char aDeviceName[256]; - anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_NAME, sizeof(aDeviceName), aDeviceName, NULL); - theInfo.Bind ("Name", aDeviceName); - - char aDeviceVendor[256]; - anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_VENDOR, sizeof(aDeviceVendor), aDeviceVendor, NULL); - theInfo.Bind ("Vendor", aDeviceVendor); - - cl_device_type aDeviceType; - anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_TYPE, sizeof(aDeviceType), &aDeviceType, NULL); - theInfo.Bind ("Type", aDeviceType == CL_DEVICE_TYPE_GPU ? "GPU" : "CPU"); - return Standard_True; -} - -// ======================================================================= -// function : ReleaseOpenCL -// purpose : Releases resources of OpenCL objects -// ======================================================================= -void OpenGl_Workspace::ReleaseOpenCL() -{ - clReleaseKernel (myRaytraceRenderKernel); - clReleaseKernel (myRaytraceSmoothKernel); - - clReleaseProgram (myRaytraceProgram); - clReleaseCommandQueue (myComputeQueue); - - clReleaseMemObject (myRaytraceOutputImage); - clReleaseMemObject (myRaytraceEnvironment); - clReleaseMemObject (myRaytraceOutputImageAA); - - clReleaseMemObject (myRaytraceMaterialBuffer); - clReleaseMemObject (myRaytraceLightSourceBuffer); - - clReleaseMemObject (mySceneNodeInfoBuffer); - clReleaseMemObject (mySceneMinPointBuffer); - clReleaseMemObject (mySceneMaxPointBuffer); - - clReleaseMemObject (myObjectNodeInfoBuffer); - clReleaseMemObject (myObjectMinPointBuffer); - clReleaseMemObject (myObjectMaxPointBuffer); - - clReleaseMemObject (myGeometryVertexBuffer); - clReleaseMemObject (myGeometryNormalBuffer); - clReleaseMemObject (myGeometryTriangBuffer); - - clReleaseContext (myComputeContext); - - if (!myGlContext.IsNull()) - { - if (!myRaytraceOutputTexture.IsNull()) - myGlContext->DelayedRelease (myRaytraceOutputTexture); - myRaytraceOutputTexture.Nullify(); - - if (!myRaytraceOutputTextureAA.IsNull()) - myGlContext->DelayedRelease (myRaytraceOutputTextureAA); - myRaytraceOutputTextureAA.Nullify(); - } -} - -// ======================================================================= -// function : ResizeRaytraceOutputBuffer -// purpose : Resizes OpenCL output image -// ======================================================================= -Standard_Boolean OpenGl_Workspace::ResizeRaytraceOutputBuffer (const cl_int theSizeX, - const cl_int theSizeY) -{ - if (myComputeContext == NULL) - { - return Standard_False; - } - - if (!myRaytraceOutputTexture.IsNull()) - { - Standard_Boolean toResize = myRaytraceOutputTexture->SizeX() != theSizeX || - myRaytraceOutputTexture->SizeY() != theSizeY; - - if (!toResize) - return Standard_True; - - if (!myGlContext.IsNull()) + if (!myGlContext->IsGlGreaterEqual (3, 1)) { - if (!myRaytraceOutputTexture.IsNull()) - myGlContext->DelayedRelease (myRaytraceOutputTexture); - if (!myRaytraceOutputTextureAA.IsNull()) - myGlContext->DelayedRelease (myRaytraceOutputTextureAA); + const TCollection_ExtendedString aMessage = "Ray-tracing requires OpenGL 3.1 and higher"; + + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); + + return Standard_False; + } + + TCollection_AsciiString aFolder = Graphic3d_ShaderProgram::ShadersFolder(); + + if (aFolder.IsEmpty()) + { + const TCollection_ExtendedString aMessage = "Failed to locate shaders directory"; + + myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, + GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); + + return Standard_False; + } + + if (myIsRaytraceDataValid) + { + myTraversalStackSize = Max (THE_DEFAULT_STACK_SIZE, + myRaytraceGeometry.HighLevelTreeDepth() + myRaytraceGeometry.BottomLevelTreeDepth()); + } + + { + Handle(OpenGl_ShaderObject) aBasicVertShader = LoadShader ( + ShaderSource (aFolder + "/RaytraceBase.vs"), GL_VERTEX_SHADER); + + if (aBasicVertShader.IsNull()) + { + return SafeFailBack ("Failed to set vertex shader source"); + } + + TCollection_AsciiString aFiles[] = { aFolder + "/RaytraceBase.fs", aFolder + "/RaytraceRender.fs" }; + + myRaytraceShaderSource.Load (aFiles, 2); + + TCollection_AsciiString aStackSizeStr = + TCollection_AsciiString ("#define STACK_SIZE ") + TCollection_AsciiString (myTraversalStackSize); + + myRaytraceShaderSource.SetPrefix (aStackSizeStr); + + myRaytraceShader = LoadShader (myRaytraceShaderSource, GL_FRAGMENT_SHADER); + + if (myRaytraceShader.IsNull()) + { + aBasicVertShader->Release (myGlContext.operator->()); + + return SafeFailBack ("Failed to set ray-trace fragment shader source"); + } + + myRaytraceProgram = new OpenGl_ShaderProgram; + + if (!myRaytraceProgram->Create (myGlContext)) + { + aBasicVertShader->Release (myGlContext.operator->()); + + return SafeFailBack ("Failed to create ray-trace shader program"); + } + + if (!myRaytraceProgram->AttachShader (myGlContext, aBasicVertShader) + || !myRaytraceProgram->AttachShader (myGlContext, myRaytraceShader)) + { + aBasicVertShader->Release (myGlContext.operator->()); + + return SafeFailBack ("Failed to attach ray-trace shader objects"); + } + + if (!myRaytraceProgram->Link (myGlContext)) + { + TCollection_AsciiString aLinkLog; + + if (myRaytraceProgram->FetchInfoLog (myGlContext, aLinkLog)) + { + #ifdef RAY_TRACE_PRINT_INFO + std::cout << aLinkLog << std::endl; + #endif + } + + return SafeFailBack ("Failed to link ray-trace shader program"); + } + } + + { + Handle(OpenGl_ShaderObject) aBasicVertShader = LoadShader ( + ShaderSource (aFolder + "/RaytraceBase.vs"), GL_VERTEX_SHADER); + + if (aBasicVertShader.IsNull()) + { + return SafeFailBack ("Failed to set vertex shader source"); + } + + TCollection_AsciiString aFiles[] = { aFolder + "/RaytraceBase.fs", aFolder + "/RaytraceSmooth.fs" }; + + myPostFSAAShaderSource.Load (aFiles, 2); + + TCollection_AsciiString aStackSizeStr = + TCollection_AsciiString ("#define STACK_SIZE ") + TCollection_AsciiString (myTraversalStackSize); + + myPostFSAAShaderSource.SetPrefix (aStackSizeStr); + + myPostFSAAShader = LoadShader (myPostFSAAShaderSource, GL_FRAGMENT_SHADER); + + if (myPostFSAAShader.IsNull()) + { + aBasicVertShader->Release (myGlContext.operator->()); + + return SafeFailBack ("Failed to set FSAA fragment shader source"); + } + + myPostFSAAProgram = new OpenGl_ShaderProgram; + + if (!myPostFSAAProgram->Create (myGlContext)) + { + aBasicVertShader->Release (myGlContext.operator->()); + + return SafeFailBack ("Failed to create FSAA shader program"); + } + + if (!myPostFSAAProgram->AttachShader (myGlContext, aBasicVertShader) + || !myPostFSAAProgram->AttachShader (myGlContext, myPostFSAAShader)) + { + aBasicVertShader->Release (myGlContext.operator->()); + + return SafeFailBack ("Failed to attach FSAA shader objects"); + } + + if (!myPostFSAAProgram->Link (myGlContext)) + { + TCollection_AsciiString aLinkLog; + + if (myPostFSAAProgram->FetchInfoLog (myGlContext, aLinkLog)) + { + #ifdef RAY_TRACE_PRINT_INFO + std::cout << aLinkLog << std::endl; + #endif + } + + return SafeFailBack ("Failed to link FSAA shader program"); + } } } - myRaytraceOutputTexture = new OpenGl_Texture(); - - myRaytraceOutputTexture->Create (myGlContext); - myRaytraceOutputTexture->InitRectangle (myGlContext, - theSizeX, theSizeY, OpenGl_TextureFormat::Create()); - - myRaytraceOutputTextureAA = new OpenGl_Texture(); - - myRaytraceOutputTextureAA->Create (myGlContext); - myRaytraceOutputTextureAA->InitRectangle (myGlContext, - theSizeX, theSizeY, OpenGl_TextureFormat::Create()); - - if (myRaytraceOutputImage != NULL) - clReleaseMemObject (myRaytraceOutputImage); - - if (myRaytraceOutputImageAA != NULL) - clReleaseMemObject (myRaytraceOutputImageAA); - - cl_int anError = CL_SUCCESS; - - myRaytraceOutputImage = clCreateFromGLTexture2D (myComputeContext, - CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTexture->TextureId(), &anError); - - if (anError != CL_SUCCESS) + if (myComputeInitStatus == OpenGl_RT_NONE || aToRebuildShaders) { -#ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to create output image!" << std::endl; -#endif - return Standard_False; + for (Standard_Integer anIndex = 0; anIndex < 2; ++anIndex) + { + Handle(OpenGl_ShaderProgram)& aShaderProgram = + (anIndex == 0) ? myRaytraceProgram : myPostFSAAProgram; + + aShaderProgram->Bind (myGlContext); + + aShaderProgram->SetSampler (myGlContext, + "uSceneMinPointTexture", OpenGl_RT_SceneMinPointTexture); + aShaderProgram->SetSampler (myGlContext, + "uSceneMaxPointTexture", OpenGl_RT_SceneMaxPointTexture); + aShaderProgram->SetSampler (myGlContext, + "uSceneNodeInfoTexture", OpenGl_RT_SceneNodeInfoTexture); + aShaderProgram->SetSampler (myGlContext, + "uObjectMinPointTexture", OpenGl_RT_ObjectMinPointTexture); + aShaderProgram->SetSampler (myGlContext, + "uObjectMaxPointTexture", OpenGl_RT_ObjectMaxPointTexture); + aShaderProgram->SetSampler (myGlContext, + "uObjectNodeInfoTexture", OpenGl_RT_ObjectNodeInfoTexture); + aShaderProgram->SetSampler (myGlContext, + "uGeometryVertexTexture", OpenGl_RT_GeometryVertexTexture); + aShaderProgram->SetSampler (myGlContext, + "uGeometryNormalTexture", OpenGl_RT_GeometryNormalTexture); + aShaderProgram->SetSampler (myGlContext, + "uGeometryTriangTexture", OpenGl_RT_GeometryTriangTexture); + aShaderProgram->SetSampler (myGlContext, + "uRaytraceMaterialTexture", OpenGl_RT_RaytraceMaterialTexture); + aShaderProgram->SetSampler (myGlContext, + "uRaytraceLightSrcTexture", OpenGl_RT_RaytraceLightSrcTexture); + + if (anIndex == 1) + { + aShaderProgram->SetSampler (myGlContext, + "uFSAAInputTexture", OpenGl_RT_FSAAInputTexture); + } + + myUniformLocations[anIndex][OpenGl_RT_aPosition] = + aShaderProgram->GetAttributeLocation (myGlContext, "aPosition"); + + myUniformLocations[anIndex][OpenGl_RT_uOriginLB] = + aShaderProgram->GetUniformLocation (myGlContext, "uOriginLB"); + myUniformLocations[anIndex][OpenGl_RT_uOriginRB] = + aShaderProgram->GetUniformLocation (myGlContext, "uOriginRB"); + myUniformLocations[anIndex][OpenGl_RT_uOriginLT] = + aShaderProgram->GetUniformLocation (myGlContext, "uOriginLT"); + myUniformLocations[anIndex][OpenGl_RT_uOriginRT] = + aShaderProgram->GetUniformLocation (myGlContext, "uOriginRT"); + myUniformLocations[anIndex][OpenGl_RT_uDirectLB] = + aShaderProgram->GetUniformLocation (myGlContext, "uDirectLB"); + myUniformLocations[anIndex][OpenGl_RT_uDirectRB] = + aShaderProgram->GetUniformLocation (myGlContext, "uDirectRB"); + myUniformLocations[anIndex][OpenGl_RT_uDirectLT] = + aShaderProgram->GetUniformLocation (myGlContext, "uDirectLT"); + myUniformLocations[anIndex][OpenGl_RT_uDirectRT] = + aShaderProgram->GetUniformLocation (myGlContext, "uDirectRT"); + + myUniformLocations[anIndex][OpenGl_RT_uLightCount] = + aShaderProgram->GetUniformLocation (myGlContext, "uLightCount"); + myUniformLocations[anIndex][OpenGl_RT_uLightAmbnt] = + aShaderProgram->GetUniformLocation (myGlContext, "uGlobalAmbient"); + + myUniformLocations[anIndex][OpenGl_RT_uSceneRad] = + aShaderProgram->GetUniformLocation (myGlContext, "uSceneRadius"); + myUniformLocations[anIndex][OpenGl_RT_uSceneEps] = + aShaderProgram->GetUniformLocation (myGlContext, "uSceneEpsilon"); + + myUniformLocations[anIndex][OpenGl_RT_uShadEnabled] = + aShaderProgram->GetUniformLocation (myGlContext, "uShadowsEnable"); + myUniformLocations[anIndex][OpenGl_RT_uReflEnabled] = + aShaderProgram->GetUniformLocation (myGlContext, "uReflectionsEnable"); + + myUniformLocations[anIndex][OpenGl_RT_uOffsetX] = + aShaderProgram->GetUniformLocation (myGlContext, "uOffsetX"); + myUniformLocations[anIndex][OpenGl_RT_uOffsetY] = + aShaderProgram->GetUniformLocation (myGlContext, "uOffsetY"); + myUniformLocations[anIndex][OpenGl_RT_uSamples] = + aShaderProgram->GetUniformLocation (myGlContext, "uSamples"); + } + + OpenGl_ShaderProgram::Unbind (myGlContext); } - myRaytraceOutputImageAA = clCreateFromGLTexture2D (myComputeContext, - CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTextureAA->TextureId(), &anError); - - if (anError != CL_SUCCESS) + if (myComputeInitStatus != OpenGl_RT_NONE) { -#ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to create anti-aliased output image!" << std::endl; -#endif - return Standard_False; + return myComputeInitStatus == OpenGl_RT_INIT; } + if (myRaytraceFBO1.IsNull()) + { + myRaytraceFBO1 = new OpenGl_FrameBuffer; + } + + if (myRaytraceFBO2.IsNull()) + { + myRaytraceFBO2 = new OpenGl_FrameBuffer; + } + + const GLfloat aVertices[] = { -1.f, -1.f, 0.f, + -1.f, 1.f, 0.f, + 1.f, 1.f, 0.f, + 1.f, 1.f, 0.f, + 1.f, -1.f, 0.f, + -1.f, -1.f, 0.f }; + + myRaytraceScreenQuad.Init (myGlContext, 3, 6, aVertices); + + myComputeInitStatus = OpenGl_RT_INIT; // initialized in normal way + return Standard_True; } // ======================================================================= -// function : WriteRaytraceSceneToDevice -// purpose : Writes scene geometry to OpenCL device +// function : NullifyResource +// purpose : // ======================================================================= -Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice() +inline void NullifyResource (const Handle(OpenGl_Context)& theContext, + Handle(OpenGl_Resource)& theResource) { - if (myComputeContext == NULL) - return Standard_False; + if (!theResource.IsNull()) + { + theResource->Release (theContext.operator->()); + theResource.Nullify(); + } +} - cl_int anErrorRes = CL_SUCCESS; +// ======================================================================= +// function : ReleaseRaytraceResources +// purpose : Releases OpenGL/GLSL shader programs +// ======================================================================= +void OpenGl_Workspace::ReleaseRaytraceResources() +{ + NullifyResource (myGlContext, myRaytraceFBO1); + NullifyResource (myGlContext, myRaytraceFBO2); - if (mySceneNodeInfoBuffer != NULL) - anErrorRes |= clReleaseMemObject (mySceneNodeInfoBuffer); + NullifyResource (myGlContext, myRaytraceShader); + NullifyResource (myGlContext, myPostFSAAShader); - if (mySceneMinPointBuffer != NULL) - anErrorRes |= clReleaseMemObject (mySceneMinPointBuffer); + NullifyResource (myGlContext, myRaytraceProgram); + NullifyResource (myGlContext, myPostFSAAProgram); - if (mySceneMaxPointBuffer != NULL) - anErrorRes |= clReleaseMemObject (mySceneMaxPointBuffer); + NullifyResource (myGlContext, mySceneNodeInfoTexture); + NullifyResource (myGlContext, mySceneMinPointTexture); + NullifyResource (myGlContext, mySceneMaxPointTexture); - if (myObjectNodeInfoBuffer != NULL) - anErrorRes |= clReleaseMemObject (myObjectNodeInfoBuffer); + NullifyResource (myGlContext, myObjectNodeInfoTexture); + NullifyResource (myGlContext, myObjectMinPointTexture); + NullifyResource (myGlContext, myObjectMaxPointTexture); - if (myObjectMinPointBuffer != NULL) - anErrorRes |= clReleaseMemObject (myObjectMinPointBuffer); + NullifyResource (myGlContext, myGeometryVertexTexture); + NullifyResource (myGlContext, myGeometryNormalTexture); + NullifyResource (myGlContext, myGeometryTriangTexture); - if (myObjectMaxPointBuffer != NULL) - anErrorRes |= clReleaseMemObject (myObjectMaxPointBuffer); + NullifyResource (myGlContext, myRaytraceLightSrcTexture); + NullifyResource (myGlContext, myRaytraceMaterialTexture); - if (myGeometryVertexBuffer != NULL) - anErrorRes |= clReleaseMemObject (myGeometryVertexBuffer); + if (myRaytraceScreenQuad.IsValid()) + myRaytraceScreenQuad.Release (myGlContext.operator->()); +} - if (myGeometryNormalBuffer != NULL) - anErrorRes |= clReleaseMemObject (myGeometryNormalBuffer); - - if (myGeometryTriangBuffer != NULL) - anErrorRes |= clReleaseMemObject (myGeometryTriangBuffer); - - if (myRaytraceMaterialBuffer != NULL) - anErrorRes |= clReleaseMemObject (myRaytraceMaterialBuffer); - - if (anErrorRes != CL_SUCCESS) +// ======================================================================= +// function : UploadRaytraceData +// purpose : Uploads ray-trace data to the GPU +// ======================================================================= +Standard_Boolean OpenGl_Workspace::UploadRaytraceData() +{ + if (!myGlContext->IsGlGreaterEqual (3, 1)) { #ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to release OpenCL buffers" << std::endl; + std::cout << "Error: OpenGL version is less than 3.1" << std::endl; #endif return Standard_False; } ///////////////////////////////////////////////////////////////////////////// - // Create material buffer + // Create OpenGL texture buffers - const size_t aMaterialBufferSize = - myRaytraceGeometry.Materials.size() != 0 ? myRaytraceGeometry.Materials.size() : 1; - - myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL, &anErrorRes); - - if (anErrorRes != CL_SUCCESS) + if (mySceneNodeInfoTexture.IsNull()) // create hight-level BVH buffers { + mySceneNodeInfoTexture = new OpenGl_TextureBufferArb; + mySceneMinPointTexture = new OpenGl_TextureBufferArb; + mySceneMaxPointTexture = new OpenGl_TextureBufferArb; + + if (!mySceneNodeInfoTexture->Create (myGlContext) + || !mySceneMinPointTexture->Create (myGlContext) + || !mySceneMaxPointTexture->Create (myGlContext)) + { #ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to create OpenCL material buffer" << std::endl; + std::cout << "Error: Failed to create buffers for high-level scene BVH" << std::endl; #endif - return Standard_False; + return Standard_False; + } + } + + if (myObjectNodeInfoTexture.IsNull()) // create bottom-level BVH buffers + { + myObjectNodeInfoTexture = new OpenGl_TextureBufferArb; + myObjectMinPointTexture = new OpenGl_TextureBufferArb; + myObjectMaxPointTexture = new OpenGl_TextureBufferArb; + + if (!myObjectNodeInfoTexture->Create (myGlContext) + || !myObjectMinPointTexture->Create (myGlContext) + || !myObjectMaxPointTexture->Create (myGlContext)) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error: Failed to create buffers for bottom-level scene BVH" << std::endl; +#endif + return Standard_False; + } + } + + if (myGeometryVertexTexture.IsNull()) // create geometry buffers + { + myGeometryVertexTexture = new OpenGl_TextureBufferArb; + myGeometryNormalTexture = new OpenGl_TextureBufferArb; + myGeometryTriangTexture = new OpenGl_TextureBufferArb; + + if (!myGeometryVertexTexture->Create (myGlContext) + || !myGeometryNormalTexture->Create (myGlContext) + || !myGeometryTriangTexture->Create (myGlContext)) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error: Failed to create buffers for triangulation data" << std::endl; +#endif + return Standard_False; + } + } + + if (myRaytraceMaterialTexture.IsNull()) // create material buffer + { + myRaytraceMaterialTexture = new OpenGl_TextureBufferArb; + + if (!myRaytraceMaterialTexture->Create (myGlContext)) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error: Failed to create buffers for material data" << std::endl; +#endif + return Standard_False; + } } ///////////////////////////////////////////////////////////////////////////// - // Create BVHs buffers - - cl_int anErrorTmp = CL_SUCCESS; + // Write OpenGL texture buffers const NCollection_Handle >& aBVH = myRaytraceGeometry.BVH(); - const size_t aSceneMinPointBufferSize = - aBVH->MinPointBuffer().size() != 0 ? aBVH->MinPointBuffer().size() : 1; + bool aResult = true; - mySceneMinPointBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aSceneMinPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp); - anErrorRes |= anErrorTmp; + if (!aBVH->NodeInfoBuffer().empty()) + { + aResult &= mySceneNodeInfoTexture->Init (myGlContext, 4, + aBVH->NodeInfoBuffer().size(), reinterpret_cast (&aBVH->NodeInfoBuffer().front())); - const size_t aSceneMaxPointBufferSize = - aBVH->MaxPointBuffer().size() != 0 ? aBVH->MaxPointBuffer().size() : 1; + aResult &= mySceneMinPointTexture->Init (myGlContext, 4, + aBVH->MinPointBuffer().size(), reinterpret_cast (&aBVH->MinPointBuffer().front())); - mySceneMaxPointBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aSceneMaxPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp); - anErrorRes |= anErrorTmp; + aResult &= mySceneMaxPointTexture->Init (myGlContext, 4, + aBVH->MaxPointBuffer().size(), reinterpret_cast (&aBVH->MaxPointBuffer().front())); + } - const size_t aSceneNodeInfoBufferSize = - aBVH->NodeInfoBuffer().size() != 0 ? aBVH->NodeInfoBuffer().size() : 1; - - mySceneNodeInfoBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aSceneNodeInfoBufferSize * sizeof(cl_int4), NULL, &anErrorTmp); - anErrorRes |= anErrorTmp; - - if (anErrorRes != CL_SUCCESS) + if (!aResult) { #ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to create OpenCL buffers for high-level scene BVH" << std::endl; + std::cout << "Error: Failed to upload buffers for high-level scene BVH" << std::endl; #endif return Standard_False; } - Standard_Integer aTotalVerticesNb = 0; - Standard_Integer aTotalElementsNb = 0; - Standard_Integer aTotalBVHNodesNb = 0; + Standard_Size aTotalVerticesNb = 0; + Standard_Size aTotalElementsNb = 0; + Standard_Size aTotalBVHNodesNb = 0; for (Standard_Integer anElemIndex = 0; anElemIndex < myRaytraceGeometry.Size(); ++anElemIndex) { @@ -1456,189 +1508,148 @@ Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice() myRaytraceGeometry.Objects().ChangeValue (anElemIndex).operator->()); Standard_ASSERT_RETURN (aTriangleSet != NULL, - "Error! Failed to get triangulation of OpenGL element", Standard_False); + "Error: Failed to get triangulation of OpenGL element", Standard_False); - aTotalVerticesNb += (int)aTriangleSet->Vertices.size(); - aTotalElementsNb += (int)aTriangleSet->Elements.size(); + aTotalVerticesNb += aTriangleSet->Vertices.size(); + aTotalElementsNb += aTriangleSet->Elements.size(); Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(), - "Error! Failed to get bottom-level BVH of OpenGL element", Standard_False); + "Error: Failed to get bottom-level BVH of OpenGL element", Standard_False); - aTotalBVHNodesNb += (int)aTriangleSet->BVH()->NodeInfoBuffer().size(); + aTotalBVHNodesNb += aTriangleSet->BVH()->NodeInfoBuffer().size(); } - aTotalBVHNodesNb = aTotalBVHNodesNb > 0 ? aTotalBVHNodesNb : 1; + if (aTotalBVHNodesNb != 0) + { + aResult &= myObjectNodeInfoTexture->Init ( + myGlContext, 4, aTotalBVHNodesNb, static_cast (NULL)); - myObjectNodeInfoBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_int4), NULL, &anErrorTmp); - anErrorRes |= anErrorTmp; + aResult &= myObjectMinPointTexture->Init ( + myGlContext, 4, aTotalBVHNodesNb, static_cast (NULL)); - myObjectMinPointBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp); - anErrorRes |= anErrorTmp; + aResult &= myObjectMaxPointTexture->Init ( + myGlContext, 4, aTotalBVHNodesNb, static_cast (NULL)); + } - myObjectMaxPointBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp); - anErrorRes |= anErrorTmp; - - if (anErrorRes != CL_SUCCESS) + if (!aResult) { #ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to create OpenCL buffers for bottom-level scene BVHs" << std::endl; + std::cout << "Error: Failed to upload buffers for bottom-level scene BVH" << std::endl; #endif return Standard_False; } - ///////////////////////////////////////////////////////////////////////////// - // Create geometry buffers + if (aTotalElementsNb != 0) + { + aResult &= myGeometryTriangTexture->Init ( + myGlContext, 4, aTotalElementsNb, static_cast (NULL)); + } - aTotalVerticesNb = aTotalVerticesNb > 0 ? aTotalVerticesNb : 1; + if (aTotalVerticesNb != 0) + { + aResult &= myGeometryVertexTexture->Init ( + myGlContext, 4, aTotalVerticesNb, static_cast (NULL)); - myGeometryVertexBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp); - anErrorRes |= anErrorTmp; + aResult &= myGeometryNormalTexture->Init ( + myGlContext, 4, aTotalVerticesNb, static_cast (NULL)); + } - myGeometryNormalBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp); - anErrorRes |= anErrorTmp; - - aTotalElementsNb = aTotalElementsNb > 0 ? aTotalElementsNb : 1; - - myGeometryTriangBuffer = clCreateBuffer (myComputeContext, - CL_MEM_READ_ONLY, aTotalElementsNb * sizeof(cl_int4), NULL, &anErrorTmp); - anErrorRes |= anErrorTmp; - - if (anErrorRes != CL_SUCCESS) + if (!aResult) { #ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to create OpenCL geometry buffers" << std::endl; + std::cout << "Error: Failed to upload buffers for scene geometry" << std::endl; #endif return Standard_False; } - ///////////////////////////////////////////////////////////////////////////// - // Write BVH and geometry buffers - - if (aBVH->NodeInfoBuffer().size() != 0) + for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx) { - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneNodeInfoBuffer, CL_FALSE, 0, - aSceneNodeInfoBufferSize * sizeof(cl_int4), &aBVH->NodeInfoBuffer().front(), 0, NULL, NULL); + if (!aBVH->IsOuter (aNodeIdx)) + continue; - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMinPointBuffer, CL_FALSE, 0, - aSceneMinPointBufferSize * sizeof(cl_float4), &aBVH->MinPointBuffer().front(), 0, NULL, NULL); + OpenGl_TriangleSet* aTriangleSet = myRaytraceGeometry.TriangleSet (aNodeIdx); - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMaxPointBuffer, CL_FALSE, 0, - aSceneMaxPointBufferSize * sizeof(cl_float4), &aBVH->MaxPointBuffer().front(), 0, NULL, NULL); + Standard_ASSERT_RETURN (aTriangleSet != NULL, + "Error: Failed to get triangulation of OpenGL element", Standard_False); - anErrorRes |= clFinish (myComputeQueue); + const Standard_Integer aBVHOffset = myRaytraceGeometry.AccelerationOffset (aNodeIdx); - if (anErrorRes != CL_SUCCESS) + Standard_ASSERT_RETURN (aBVHOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET, + "Error: Failed to get offset for bottom-level BVH", Standard_False); + + const size_t aBVHBuffserSize = aTriangleSet->BVH()->NodeInfoBuffer().size(); + + if (aBVHBuffserSize != 0) + { + aResult &= myObjectNodeInfoTexture->SubData (myGlContext, aBVHOffset, + aBVHBuffserSize, reinterpret_cast (&aTriangleSet->BVH()->NodeInfoBuffer().front())); + + aResult &= myObjectMinPointTexture->SubData (myGlContext, aBVHOffset, + aBVHBuffserSize, reinterpret_cast (&aTriangleSet->BVH()->MinPointBuffer().front())); + + aResult &= myObjectMaxPointTexture->SubData (myGlContext, aBVHOffset, + aBVHBuffserSize, reinterpret_cast (&aTriangleSet->BVH()->MaxPointBuffer().front())); + + if (!aResult) + { +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error: Failed to upload buffers for bottom-level scene BVHs" << std::endl; +#endif + return Standard_False; + } + } + + const Standard_Integer aVerticesOffset = myRaytraceGeometry.VerticesOffset (aNodeIdx); + + Standard_ASSERT_RETURN (aVerticesOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET, + "Error: Failed to get offset for triangulation vertices of OpenGL element", Standard_False); + + if (!aTriangleSet->Vertices.empty()) + { + aResult &= myGeometryNormalTexture->SubData (myGlContext, aVerticesOffset, + aTriangleSet->Normals.size(), reinterpret_cast (&aTriangleSet->Normals.front())); + + aResult &= myGeometryVertexTexture->SubData (myGlContext, aVerticesOffset, + aTriangleSet->Vertices.size(), reinterpret_cast (&aTriangleSet->Vertices.front())); + } + + const Standard_Integer anElementsOffset = myRaytraceGeometry.ElementsOffset (aNodeIdx); + + Standard_ASSERT_RETURN (anElementsOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET, + "Error: Failed to get offset for triangulation elements of OpenGL element", Standard_False); + + if (!aTriangleSet->Elements.empty()) + { + aResult &= myGeometryTriangTexture->SubData (myGlContext, anElementsOffset, + aTriangleSet->Elements.size(), reinterpret_cast (&aTriangleSet->Elements.front())); + } + + if (!aResult) { #ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to write OpenCL buffers for high-level scene BVH" << std::endl; + std::cout << "Error: Failed to upload triangulation buffers for OpenGL element" << std::endl; #endif return Standard_False; } - - for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx) - { - if (!aBVH->IsOuter (aNodeIdx)) - continue; - - OpenGl_TriangleSet* aTriangleSet = myRaytraceGeometry.TriangleSet (aNodeIdx); - - Standard_ASSERT_RETURN (aTriangleSet != NULL, - "Error! Failed to get triangulation of OpenGL element", Standard_False); - - const size_t aBVHBuffserSize = - aTriangleSet->BVH()->NodeInfoBuffer().size() != 0 ? aTriangleSet->BVH()->NodeInfoBuffer().size() : 1; - - const Standard_Integer aBVHOffset = myRaytraceGeometry.AccelerationOffset (aNodeIdx); - - Standard_ASSERT_RETURN (aBVHOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET, - "Error! Failed to get offset for bottom-level BVH", Standard_False); - - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectNodeInfoBuffer, CL_FALSE, aBVHOffset * sizeof(cl_int4), - aBVHBuffserSize * sizeof(cl_int4), &aTriangleSet->BVH()->NodeInfoBuffer().front(), 0, NULL, NULL); - - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMinPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4), - aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MinPointBuffer().front(), 0, NULL, NULL); - - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMaxPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4), - aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MaxPointBuffer().front(), 0, NULL, NULL); - - anErrorRes |= clFinish (myComputeQueue); - - if (anErrorRes != CL_SUCCESS) - { -#ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to write OpenCL buffers for bottom-level scene BVHs" << std::endl; -#endif - return Standard_False; - } - - const Standard_Integer aVerticesOffset = myRaytraceGeometry.VerticesOffset (aNodeIdx); - - Standard_ASSERT_RETURN (aVerticesOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET, - "Error! Failed to get offset for triangulation vertices of OpenGL element", Standard_False); - - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryVertexBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4), - aTriangleSet->Vertices.size() * sizeof(cl_float4), &aTriangleSet->Vertices.front(), 0, NULL, NULL); - - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryNormalBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4), - aTriangleSet->Normals.size() * sizeof(cl_float4), &aTriangleSet->Normals.front(), 0, NULL, NULL); - - const Standard_Integer anElementsOffset = myRaytraceGeometry.ElementsOffset (aNodeIdx); - - Standard_ASSERT_RETURN (anElementsOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET, - "Error! Failed to get offset for triangulation elements of OpenGL element", Standard_False); - - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryTriangBuffer, CL_FALSE, anElementsOffset * sizeof(cl_int4), - aTriangleSet->Elements.size() * sizeof(cl_int4), &aTriangleSet->Elements.front(), 0, NULL, NULL); - - anErrorRes |= clFinish (myComputeQueue); - - if (anErrorRes != CL_SUCCESS) - { -#ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to write OpenCL triangulation buffers for OpenGL element" << std::endl; -#endif - return Standard_False; - } - } } - ///////////////////////////////////////////////////////////////////////////// - // Write material buffer - if (myRaytraceGeometry.Materials.size() != 0) { - const void* aDataPtr = myRaytraceGeometry.Materials.front().Packed(); + const GLfloat* aDataPtr = myRaytraceGeometry.Materials.front().Packed(); - anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceMaterialBuffer, - CL_FALSE, 0, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr, 0, NULL, NULL); + aResult &= myRaytraceMaterialTexture->Init ( + myGlContext, 4, myRaytraceGeometry.Materials.size() * 7, aDataPtr); - if (anErrorRes != CL_SUCCESS) + if (!aResult) { - #ifdef RAY_TRACE_PRINT_INFO - std::cout << "Error! Failed to write OpenCL material buffer" << std::endl; - #endif +#ifdef RAY_TRACE_PRINT_INFO + std::cout << "Error: Failed to upload material buffer" << std::endl; +#endif return Standard_False; } } - anErrorRes |= clFinish (myComputeQueue); - - if (anErrorRes == CL_SUCCESS) - { - myIsRaytraceDataValid = myRaytraceGeometry.Objects().Size() != 0; - } -#ifdef RAY_TRACE_PRINT_INFO - else - { - std::cout << "Error! Failed to set scene data buffers" << std::endl; - } -#endif + myIsRaytraceDataValid = myRaytraceGeometry.Objects().Size() != 0; #ifdef RAY_TRACE_PRINT_INFO @@ -1675,498 +1686,321 @@ Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice() #endif - return (CL_SUCCESS == anErrorRes); + return aResult; } -// Use it to estimate the optimal size of OpenCL work group -// #define OPENCL_GROUP_SIZE_TEST - // ======================================================================= -// function : RunRaytraceOpenCLKernels -// purpose : Runs OpenCL ray-tracing kernels +// function : ResizeRaytraceBuffers +// purpose : Resizes OpenGL frame buffers // ======================================================================= -Standard_Boolean OpenGl_Workspace::RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView, - const Standard_ShortReal theOrigins[16], - const Standard_ShortReal theDirects[16], - const Standard_Integer theSizeX, - const Standard_Integer theSizeY) +Standard_Boolean OpenGl_Workspace::ResizeRaytraceBuffers (const Standard_Integer theSizeX, + const Standard_Integer theSizeY) { - if (myRaytraceRenderKernel == NULL || myComputeQueue == NULL) - return Standard_False; - - //////////////////////////////////////////////////////////////////////// - // Set kernel arguments - - cl_uint anIndex = 0; - cl_int anError = 0; - - cl_int aLightSourceBufferSize = (cl_int)myRaytraceGeometry.Sources.size(); - - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theSizeX); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theSizeY); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_float16), theOrigins); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_float16), theDirects); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceEnvironment); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImage); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneNodeInfoBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneMinPointBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneMaxPointBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectNodeInfoBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectMinPointBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectMaxPointBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryTriangBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryVertexBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryNormalBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceLightSourceBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceMaterialBuffer); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_float4), &myRaytraceGeometry.GlobalAmbient); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &aLightSourceBufferSize); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theCView.IsShadowsEnabled); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theCView.IsReflectionsEnabled); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneEpsilon); - anError |= clSetKernelArg ( - myRaytraceRenderKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneRadius); - - if (anError != CL_SUCCESS) + if (myRaytraceFBO1->GetVPSizeX() != theSizeX + || myRaytraceFBO1->GetVPSizeY() != theSizeY) { - const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of ray-tracing kernel!"; - - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); - - return Standard_False; + myRaytraceFBO1->Init (myGlContext, theSizeX, theSizeY); + myRaytraceFBO2->Init (myGlContext, theSizeX, theSizeY); } - // Second-pass 'smoothing' kernel runs only if anti-aliasing is enabled - if (theCView.IsAntialiasingEnabled) - { - anIndex = 0; - - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theSizeX); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theSizeY); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_float16), theOrigins); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_float16), theDirects); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImage); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceEnvironment); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImageAA); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneNodeInfoBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneMinPointBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneMaxPointBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectNodeInfoBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectMinPointBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectMaxPointBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryTriangBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryVertexBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryNormalBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceLightSourceBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceMaterialBuffer); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_float4), &myRaytraceGeometry.GlobalAmbient); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &aLightSourceBufferSize); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theCView.IsShadowsEnabled); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theCView.IsReflectionsEnabled); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneEpsilon); - anError |= clSetKernelArg ( - myRaytraceSmoothKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneRadius); - - if (anError != CL_SUCCESS) - { - const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of smoothing kernel!"; - - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); - - return Standard_False; - } - } - - //////////////////////////////////////////////////////////////////////// - // Set work size - - size_t aLocWorkSize[] = { myIsAmdComputePlatform ? 2 : 4, 32 }; - -#ifdef OPENCL_GROUP_SIZE_TEST - for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1) - for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1) - { - aLocWorkSize[0] = aLocX; - aLocWorkSize[1] = aLocY; -#endif - - size_t aWorkSizeX = theSizeX; - if (aWorkSizeX % aLocWorkSize[0] != 0) - aWorkSizeX += aLocWorkSize[0] - aWorkSizeX % aLocWorkSize[0]; - - size_t aWokrSizeY = theSizeY; - if (aWokrSizeY % aLocWorkSize[1] != 0 ) - aWokrSizeY += aLocWorkSize[1] - aWokrSizeY % aLocWorkSize[1]; - - size_t aTotWorkSize[] = { aWorkSizeX, aWokrSizeY }; - - cl_event anEvent = NULL, anEventSmooth = NULL; - - anError = clEnqueueNDRangeKernel (myComputeQueue, - myRaytraceRenderKernel, 2, NULL, aTotWorkSize, aLocWorkSize, 0, NULL, &anEvent); - - if (anError != CL_SUCCESS) - { - const TCollection_ExtendedString aMessage = "Error! Failed to execute the ray-tracing kernel!"; - - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); - - return Standard_False; - } - - clWaitForEvents (1, &anEvent); - - if (theCView.IsAntialiasingEnabled) - { - size_t aLocWorkSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4, - myIsAmdComputePlatform ? 8 : 32 }; - -#ifdef OPENCL_GROUP_SIZE_TEST - aLocWorkSizeSmooth[0] = aLocX; - aLocWorkSizeSmooth[1] = aLocY; -#endif - - aWorkSizeX = theSizeX; - if (aWorkSizeX % aLocWorkSizeSmooth[0] != 0) - aWorkSizeX += aLocWorkSizeSmooth[0] - aWorkSizeX % aLocWorkSizeSmooth[0]; - - size_t aWokrSizeY = theSizeY; - if (aWokrSizeY % aLocWorkSizeSmooth[1] != 0 ) - aWokrSizeY += aLocWorkSizeSmooth[1] - aWokrSizeY % aLocWorkSizeSmooth[1]; - - size_t aTotWorkSizeSmooth [] = { aWorkSizeX, aWokrSizeY }; - - anError = clEnqueueNDRangeKernel (myComputeQueue, myRaytraceSmoothKernel, - 2, NULL, aTotWorkSizeSmooth, aLocWorkSizeSmooth, 0, NULL, &anEventSmooth); - - clWaitForEvents (1, &anEventSmooth); - - if (anError != CL_SUCCESS) - { - const TCollection_ExtendedString aMessage = "Error! Failed to execute the smoothing kernel"; - - myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, - GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); - - return Standard_False; - } - } - -#if defined (RAY_TRACE_PRINT_INFO) || defined (OPENCL_GROUP_SIZE_TEST) - - static cl_ulong ttt1 = 10000000000; - static cl_ulong ttt2 = 10000000000; - - cl_ulong aBegTime = 0; - cl_ulong aEndTime = 0; - - clGetEventProfilingInfo (anEvent, - CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL); - clGetEventProfilingInfo (anEvent, - CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL); - - ttt1 = aEndTime - aBegTime < ttt1 ? aEndTime - aBegTime : ttt1; - - std::cout << "\tRender time (ms): " << ttt1 / 1e6f << std::endl; - - if (theCView.IsAntialiasingEnabled) - { - clGetEventProfilingInfo (anEventSmooth, - CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL); - clGetEventProfilingInfo (anEventSmooth, - CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL); - - ttt2 = aEndTime - aBegTime < ttt2 ? aEndTime - aBegTime : ttt2; - - std::cout << "\tSmooth time (ms): " << ttt2 / 1e6f << std::endl; - } - -#endif - - if (anEvent != NULL) - clReleaseEvent (anEvent); - - if (anEventSmooth != NULL) - clReleaseEvent (anEventSmooth); - -#ifdef OPENCL_GROUP_SIZE_TEST - } -#endif - return Standard_True; } // ======================================================================= -// function : ComputeInverseMatrix -// purpose : Computes inversion of 4x4 floating-point matrix +// function : UpdateCamera +// purpose : Generates viewing rays for corners of screen quad // ======================================================================= -template -void ComputeInverseMatrix (const T m[16], T inv[16]) +void OpenGl_Workspace::UpdateCamera (const NCollection_Mat4& theOrientation, + const NCollection_Mat4& theViewMapping, + OpenGl_Vec3 theOrigins[4], + OpenGl_Vec3 theDirects[4]) { - inv[ 0] = m[ 5] * (m[10] * m[15] - m[11] * m[14]) - - m[ 9] * (m[ 6] * m[15] - m[ 7] * m[14]) - - m[13] * (m[ 7] * m[10] - m[ 6] * m[11]); + NCollection_Mat4 aInvModelProj; - inv[ 1] = m[ 1] * (m[11] * m[14] - m[10] * m[15]) - - m[ 9] * (m[ 3] * m[14] - m[ 2] * m[15]) - - m[13] * (m[ 2] * m[11] - m[ 3] * m[10]); + // compute invserse model-view-projection matrix + (theViewMapping * theOrientation).Inverted (aInvModelProj); - inv[ 2] = m[ 1] * (m[ 6] * m[15] - m[ 7] * m[14]) - - m[ 5] * (m[ 2] * m[15] - m[ 3] * m[14]) - - m[13] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]); - - inv[ 3] = m[ 1] * (m[ 7] * m[10] - m[ 6] * m[11]) - - m[ 5] * (m[ 3] * m[10] - m[ 2] * m[11]) - - m[ 9] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]); - - inv[ 4] = m[ 4] * (m[11] * m[14] - m[10] * m[15]) - - m[ 8] * (m[ 7] * m[14] - m[ 6] * m[15]) - - m[12] * (m[ 6] * m[11] - m[ 7] * m[10]); - - inv[ 5] = m[ 0] * (m[10] * m[15] - m[11] * m[14]) - - m[ 8] * (m[ 2] * m[15] - m[ 3] * m[14]) - - m[12] * (m[ 3] * m[10] - m[ 2] * m[11]); - - inv[ 6] = m[ 0] * (m[ 7] * m[14] - m[ 6] * m[15]) - - m[ 4] * (m[ 3] * m[14] - m[ 2] * m[15]) - - m[12] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]); - - inv[ 7] = m[ 0] * (m[ 6] * m[11] - m[ 7] * m[10]) - - m[ 4] * (m[ 2] * m[11] - m[ 3] * m[10]) - - m[ 8] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]); - - inv[ 8] = m[ 4] * (m[ 9] * m[15] - m[11] * m[13]) - - m[ 8] * (m[ 5] * m[15] - m[ 7] * m[13]) - - m[12] * (m[ 7] * m[ 9] - m[ 5] * m[11]); - - inv[ 9] = m[ 0] * (m[11] * m[13] - m[ 9] * m[15]) - - m[ 8] * (m[ 3] * m[13] - m[ 1] * m[15]) - - m[12] * (m[ 1] * m[11] - m[ 3] * m[ 9]); - - inv[10] = m[ 0] * (m[ 5] * m[15] - m[ 7] * m[13]) - - m[ 4] * (m[ 1] * m[15] - m[ 3] * m[13]) - - m[12] * (m[ 3] * m[ 5] - m[ 1] * m[ 7]); - - inv[11] = m[ 0] * (m[ 7] * m[ 9] - m[ 5] * m[11]) - - m[ 4] * (m[ 3] * m[ 9] - m[ 1] * m[11]) - - m[ 8] * (m[ 1] * m[ 7] - m[ 3] * m[ 5]); - - inv[12] = m[ 4] * (m[10] * m[13] - m[ 9] * m[14]) - - m[ 8] * (m[ 6] * m[13] - m[ 5] * m[14]) - - m[12] * (m[ 5] * m[10] - m[ 6] * m[ 9]); - - inv[13] = m[ 0] * (m[ 9] * m[14] - m[10] * m[13]) - - m[ 8] * (m[ 1] * m[14] - m[ 2] * m[13]) - - m[12] * (m[ 2] * m[ 9] - m[ 1] * m[10]); - - inv[14] = m[ 0] * (m[ 6] * m[13] - m[ 5] * m[14]) - - m[ 4] * (m[ 2] * m[13] - m[ 1] * m[14]) - - m[12] * (m[ 1] * m[ 6] - m[ 2] * m[ 5]); - - inv[15] = m[ 0] * (m[ 5] * m[10] - m[ 6] * m[ 9]) - - m[ 4] * (m[ 1] * m[10] - m[ 2] * m[ 9]) - - m[ 8] * (m[ 2] * m[ 5] - m[ 1] * m[ 6]); - - T det = m[0] * inv[ 0] + - m[1] * inv[ 4] + - m[2] * inv[ 8] + - m[3] * inv[12]; - - if (det == T (0.0)) return; - - det = T (1.0) / det; - - for (Standard_Integer i = 0; i < 16; ++i) - inv[i] *= det; -} - -// ======================================================================= -// function : GenerateCornerRays -// purpose : Generates primary rays for corners of screen quad -// ======================================================================= -void GenerateCornerRays (const GLdouble theInvModelProj[16], - cl_float theOrigins[16], - cl_float theDirects[16]) -{ Standard_Integer aOriginIndex = 0; Standard_Integer aDirectIndex = 0; - for (Standard_Integer y = -1; y <= 1; y += 2) + for (Standard_Integer aY = -1; aY <= 1; aY += 2) { - for (Standard_Integer x = -1; x <= 1; x += 2) + for (Standard_Integer aX = -1; aX <= 1; aX += 2) { - BVH_Vec4f aOrigin (float(x), - float(y), - -1.f, - 1.f); + OpenGl_Vec4d aOrigin (GLdouble(aX), + GLdouble(aY), + -1.0, + 1.0); + + aOrigin = aInvModelProj * aOrigin; - aOrigin = MatVecMult (theInvModelProj, aOrigin); aOrigin.x() = aOrigin.x() / aOrigin.w(); aOrigin.y() = aOrigin.y() / aOrigin.w(); aOrigin.z() = aOrigin.z() / aOrigin.w(); - aOrigin.w() = 1.f; - BVH_Vec4f aDirect (float(x), - float(y), - 1.f, - 1.f); + OpenGl_Vec4d aDirect (GLdouble(aX), + GLdouble(aY), + 1.0, + 1.0); + + aDirect = aInvModelProj * aDirect; - aDirect = MatVecMult (theInvModelProj, aDirect); aDirect.x() = aDirect.x() / aDirect.w(); aDirect.y() = aDirect.y() / aDirect.w(); aDirect.z() = aDirect.z() / aDirect.w(); - aDirect.w() = 1.f; aDirect = aDirect - aOrigin; - GLdouble aInvLen = 1.f / sqrt (aDirect.x() * aDirect.x() + + GLdouble aInvLen = 1.0 / sqrt (aDirect.x() * aDirect.x() + aDirect.y() * aDirect.y() + aDirect.z() * aDirect.z()); - theOrigins [aOriginIndex++] = static_cast (aOrigin.x()); - theOrigins [aOriginIndex++] = static_cast (aOrigin.y()); - theOrigins [aOriginIndex++] = static_cast (aOrigin.z()); - theOrigins [aOriginIndex++] = 1.f; + theOrigins[aOriginIndex++] = OpenGl_Vec3 (static_cast (aOrigin.x()), + static_cast (aOrigin.y()), + static_cast (aOrigin.z())); - theDirects [aDirectIndex++] = static_cast (aDirect.x() * aInvLen); - theDirects [aDirectIndex++] = static_cast (aDirect.y() * aInvLen); - theDirects [aDirectIndex++] = static_cast (aDirect.z() * aInvLen); - theDirects [aDirectIndex++] = 0.f; + theDirects[aDirectIndex++] = OpenGl_Vec3 (static_cast (aDirect.x() * aInvLen), + static_cast (aDirect.y() * aInvLen), + static_cast (aDirect.z() * aInvLen)); } } } +// ======================================================================= +// function : RunRaytraceShaders +// purpose : Runs ray-tracing shader programs +// ======================================================================= +Standard_Boolean OpenGl_Workspace::RunRaytraceShaders (const Graphic3d_CView& theCView, + const Standard_Integer theSizeX, + const Standard_Integer theSizeY, + const OpenGl_Vec3 theOrigins[4], + const OpenGl_Vec3 theDirects[4], + OpenGl_FrameBuffer* theFrameBuffer) +{ + mySceneMinPointTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_SceneMinPointTexture); + mySceneMaxPointTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_SceneMaxPointTexture); + mySceneNodeInfoTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_SceneNodeInfoTexture); + myObjectMinPointTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_ObjectMinPointTexture); + myObjectMaxPointTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_ObjectMaxPointTexture); + myObjectNodeInfoTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_ObjectNodeInfoTexture); + myGeometryVertexTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_GeometryVertexTexture); + myGeometryNormalTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_GeometryNormalTexture); + myGeometryTriangTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_GeometryTriangTexture); + myRaytraceMaterialTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_RaytraceMaterialTexture); + myRaytraceLightSrcTexture->BindTexture (myGlContext, GL_TEXTURE0 + OpenGl_RT_RaytraceLightSrcTexture); + + if (theCView.IsAntialiasingEnabled) // render source image to FBO + { + myRaytraceFBO1->BindBuffer (myGlContext); + + glDisable (GL_BLEND); + } + + myRaytraceProgram->Bind (myGlContext); + + Standard_Integer aLightSourceBufferSize = + static_cast (myRaytraceGeometry.Sources.size()); + + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uOriginLB], theOrigins[0]); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uOriginRB], theOrigins[1]); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uOriginLT], theOrigins[2]); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uOriginRT], theOrigins[3]); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uDirectLB], theDirects[0]); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uDirectRB], theDirects[1]); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uDirectLT], theDirects[2]); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uDirectRT], theDirects[3]); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uSceneRad], myRaytraceSceneRadius); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uSceneEps], myRaytraceSceneEpsilon); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uLightCount], aLightSourceBufferSize); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uLightAmbnt], myRaytraceGeometry.Ambient); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uShadEnabled], theCView.IsShadowsEnabled); + myRaytraceProgram->SetUniform (myGlContext, + myUniformLocations[0][OpenGl_RT_uReflEnabled], theCView.IsReflectionsEnabled); + + myGlContext->core20fwd->glEnableVertexAttribArray ( + myUniformLocations[0][OpenGl_RT_aPosition]); + { + myGlContext->core20fwd->glVertexAttribPointer ( + myUniformLocations[0][OpenGl_RT_aPosition], 3, GL_FLOAT, GL_FALSE, 0, NULL); + + myGlContext->core15fwd->glDrawArrays (GL_TRIANGLES, 0, 6); + } + myGlContext->core20fwd->glDisableVertexAttribArray ( + myUniformLocations[0][OpenGl_RT_aPosition]); + + if (!theCView.IsAntialiasingEnabled) + { + myRaytraceProgram->Unbind (myGlContext); + + return Standard_True; + } + + myGlContext->core20fwd->glActiveTexture ( + GL_TEXTURE0 + OpenGl_RT_FSAAInputTexture); // texture unit for FBO texture + + myRaytraceFBO1->BindTexture (myGlContext); + + myPostFSAAProgram->Bind (myGlContext); + + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uOriginLB], theOrigins[0]); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uOriginRB], theOrigins[1]); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uOriginLT], theOrigins[2]); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uOriginRT], theOrigins[3]); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uDirectLB], theDirects[0]); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uDirectRB], theDirects[1]); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uDirectLT], theDirects[2]); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uDirectRT], theDirects[3]); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uSceneRad], myRaytraceSceneRadius); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uSceneEps], myRaytraceSceneEpsilon); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uLightCount], aLightSourceBufferSize); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uLightAmbnt], myRaytraceGeometry.Ambient); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uShadEnabled], theCView.IsShadowsEnabled); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uReflEnabled], theCView.IsReflectionsEnabled); + + const Standard_ShortReal aMaxOffset = 0.559017f; + const Standard_ShortReal aMinOffset = 0.186339f; + + myGlContext->core20fwd->glEnableVertexAttribArray ( + myUniformLocations[1][OpenGl_RT_aPosition]); + + myGlContext->core20fwd->glVertexAttribPointer ( + myUniformLocations[1][OpenGl_RT_aPosition], 3, GL_FLOAT, GL_FALSE, 0, NULL); + + // Perform multi-pass adaptive FSAA using ping-pong technique + for (Standard_Integer anIt = 0; anIt < 4; ++anIt) + { + GLfloat aOffsetX = 1.f / theSizeX; + GLfloat aOffsetY = 1.f / theSizeY; + + if (anIt < 2) + { + aOffsetX *= anIt < 1 ? aMinOffset : -aMaxOffset; + aOffsetY *= anIt < 1 ? aMaxOffset : aMinOffset; + } + else + { + aOffsetX *= anIt > 2 ? aMaxOffset : -aMinOffset; + aOffsetY *= anIt > 2 ? -aMinOffset : -aMaxOffset; + } + + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uSamples], anIt + 2); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uOffsetX], aOffsetX); + myPostFSAAProgram->SetUniform (myGlContext, + myUniformLocations[1][OpenGl_RT_uOffsetY], aOffsetY); + + Handle(OpenGl_FrameBuffer)& aFramebuffer = anIt % 2 ? myRaytraceFBO1 : myRaytraceFBO2; + + if (anIt == 3) // disable FBO on last iteration + { + glEnable (GL_BLEND); + + if (theFrameBuffer != NULL) + theFrameBuffer->BindBuffer (myGlContext); + } + else + { + aFramebuffer->BindBuffer (myGlContext); + } + + myGlContext->core15fwd->glDrawArrays (GL_TRIANGLES, 0, 6); + + if (anIt != 3) // set input for the next pass + { + aFramebuffer->BindTexture (myGlContext); + aFramebuffer->UnbindBuffer (myGlContext); + } + } + + myGlContext->core20fwd->glDisableVertexAttribArray ( + myUniformLocations[1][OpenGl_RT_aPosition]); + + myPostFSAAProgram->Unbind (myGlContext); + + return Standard_True; +} + // ======================================================================= // function : Raytrace -// purpose : Redraws the window using OpenCL ray tracing +// purpose : Redraws the window using OpenGL/GLSL ray-tracing // ======================================================================= Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView, const Standard_Integer theSizeX, const Standard_Integer theSizeY, - const Tint theToSwap) + const Standard_Boolean theToSwap, + OpenGl_FrameBuffer* theFrameBuffer) { - if (!InitOpenCL()) + if (!UpdateRaytraceGeometry (Standard_True)) return Standard_False; - if (!ResizeRaytraceOutputBuffer (theSizeX, theSizeY)) + if (!InitRaytraceResources()) + return Standard_False; + + if (!ResizeRaytraceBuffers (theSizeX, theSizeY)) return Standard_False; if (!UpdateRaytraceEnvironmentMap()) return Standard_False; - if (!UpdateRaytraceGeometry (Standard_True)) - return Standard_False; - // Get model-view and projection matrices TColStd_Array2OfReal theOrientation (0, 3, 0, 3); TColStd_Array2OfReal theViewMapping (0, 3, 0, 3); myView->GetMatrices (theOrientation, theViewMapping); - GLdouble aOrientationMatrix[16]; - GLdouble aViewMappingMatrix[16]; - GLdouble aOrientationInvers[16]; + NCollection_Mat4 aOrientationMatrix; + NCollection_Mat4 aViewMappingMatrix; for (Standard_Integer j = 0; j < 4; ++j) + { for (Standard_Integer i = 0; i < 4; ++i) { aOrientationMatrix [4 * j + i] = theOrientation (i, j); aViewMappingMatrix [4 * j + i] = theViewMapping (i, j); } + } + + NCollection_Mat4 aInvOrientationMatrix; + aOrientationMatrix.Inverted (aInvOrientationMatrix); - ComputeInverseMatrix (aOrientationMatrix, aOrientationInvers); - - if (!UpdateRaytraceLightSources (aOrientationInvers)) + if (!UpdateRaytraceLightSources (aInvOrientationMatrix)) return Standard_False; - // Generate primary rays for corners of the screen quad - glMatrixMode (GL_MODELVIEW); + OpenGl_Vec3 aOrigins[4]; + OpenGl_Vec3 aDirects[4]; - glLoadMatrixd (aViewMappingMatrix); - glMultMatrixd (aOrientationMatrix); - - GLdouble aModelProject[16]; - GLdouble aInvModelProj[16]; - - glGetDoublev (GL_MODELVIEW_MATRIX, aModelProject); - - ComputeInverseMatrix (aModelProject, aInvModelProj); - - GLfloat aOrigins[16]; - GLfloat aDirects[16]; - - GenerateCornerRays (aInvModelProj, - aOrigins, - aDirects); - - // Compute ray-traced image using OpenCL kernel - cl_mem anImages[] = { myRaytraceOutputImage, myRaytraceOutputImageAA }; - cl_int anError = clEnqueueAcquireGLObjects (myComputeQueue, - 2, anImages, - 0, NULL, NULL); - clFinish (myComputeQueue); - - if (myIsRaytraceDataValid) - { - RunRaytraceOpenCLKernels (theCView, - aOrigins, - aDirects, - theSizeX, - theSizeY); - } - - anError |= clEnqueueReleaseGLObjects (myComputeQueue, - 2, anImages, - 0, NULL, NULL); - clFinish (myComputeQueue); + UpdateCamera (aOrientationMatrix, + aViewMappingMatrix, + aOrigins, + aDirects); // Draw background glPushAttrib (GL_ENABLE_BIT | @@ -2190,44 +2024,38 @@ Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView, glClear (GL_COLOR_BUFFER_BIT); - Handle(OpenGl_Workspace) aWorkspace (this); - myView->DrawBackground (aWorkspace); + if (theFrameBuffer != NULL) + theFrameBuffer->BindBuffer (myGlContext); - // Draw dummy quad to show result image - glEnable (GL_COLOR_MATERIAL); - glEnable (GL_BLEND); - - glDisable (GL_DEPTH_TEST); - - glBlendFunc (GL_ONE, GL_SRC_ALPHA); - - glEnable (GL_TEXTURE_RECTANGLE); + myView->DrawBackground (*this); + // Generate ray-traced image glMatrixMode (GL_PROJECTION); glLoadIdentity(); glMatrixMode (GL_MODELVIEW); glLoadIdentity(); - glColor3f (1.0f, 1.0f, 1.0f); - - if (!theCView.IsAntialiasingEnabled) - myRaytraceOutputTexture->Bind (myGlContext); - else - myRaytraceOutputTextureAA->Bind (myGlContext); + glEnable (GL_BLEND); + glBlendFunc (GL_ONE, GL_SRC_ALPHA); if (myIsRaytraceDataValid) { - glBegin (GL_QUADS); - { - glTexCoord2i ( 0, 0); glVertex2f (-1.f, -1.f); - glTexCoord2i ( 0, theSizeY); glVertex2f (-1.f, 1.f); - glTexCoord2i (theSizeX, theSizeY); glVertex2f ( 1.f, 1.f); - glTexCoord2i (theSizeX, 0); glVertex2f ( 1.f, -1.f); - } - glEnd(); + myRaytraceScreenQuad.Bind (myGlContext); + + RunRaytraceShaders (theCView, + theSizeX, + theSizeY, + aOrigins, + aDirects, + theFrameBuffer); + + myRaytraceScreenQuad.Unbind (myGlContext); } + if (theFrameBuffer != NULL) + theFrameBuffer->UnbindBuffer (myGlContext); + glPopAttrib(); // Swap the buffers @@ -2237,9 +2065,9 @@ Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView, myBackBufferRestored = Standard_False; } else + { glFlush(); + } return Standard_True; } - -#endif diff --git a/src/Shaders/RaytraceBase.fs b/src/Shaders/RaytraceBase.fs new file mode 100644 index 0000000000..1edc7a2f0a --- /dev/null +++ b/src/Shaders/RaytraceBase.fs @@ -0,0 +1,768 @@ +//! Normalized pixel coordinates. +in vec2 vPixel; + +//! Origin of viewing ray in left-top corner. +uniform vec3 uOriginLT; +//! Origin of viewing ray in left-bottom corner. +uniform vec3 uOriginLB; +//! Origin of viewing ray in right-top corner. +uniform vec3 uOriginRT; +//! Origin of viewing ray in right-bottom corner. +uniform vec3 uOriginRB; + +//! Direction of viewing ray in left-top corner. +uniform vec3 uDirectLT; +//! Direction of viewing ray in left-bottom corner. +uniform vec3 uDirectLB; +//! Direction of viewing ray in right-top corner. +uniform vec3 uDirectRT; +//! Direction of viewing ray in right-bottom corner. +uniform vec3 uDirectRB; + +//! Texture buffer of data records of high-level BVH nodes. +uniform isamplerBuffer uSceneNodeInfoTexture; +//! Texture buffer of minimum points of high-level BVH nodes. +uniform samplerBuffer uSceneMinPointTexture; +//! Texture buffer of maximum points of high-level BVH nodes. +uniform samplerBuffer uSceneMaxPointTexture; + +//! Texture buffer of data records of bottom-level BVH nodes. +uniform isamplerBuffer uObjectNodeInfoTexture; +//! Texture buffer of minimum points of bottom-level BVH nodes. +uniform samplerBuffer uObjectMinPointTexture; +//! Texture buffer of maximum points of bottom-level BVH nodes. +uniform samplerBuffer uObjectMaxPointTexture; + +//! Texture buffer of vertex coords. +uniform samplerBuffer uGeometryVertexTexture; +//! Texture buffer of vertex normals. +uniform samplerBuffer uGeometryNormalTexture; +//! Texture buffer of triangle indices. +uniform isamplerBuffer uGeometryTriangTexture; + +//! Texture buffer of material properties. +uniform samplerBuffer uRaytraceMaterialTexture; +//! Texture buffer of light source properties. +uniform samplerBuffer uRaytraceLightSrcTexture; +//! Environment map texture. +uniform sampler2D uEnvironmentMapTexture; + +//! Total number of light sources. +uniform int uLightCount; +//! Intensity of global ambient light. +uniform vec4 uGlobalAmbient; + +//! Enables/disables environment map. +uniform int uEnvironmentEnable; +//! Enables/disables computation of shadows. +uniform int uShadowsEnable; +//! Enables/disables computation of reflections. +uniform int uReflectionsEnable; + +//! Radius of bounding sphere of the scene. +uniform float uSceneRadius; +//! Scene epsilon to prevent self-intersections. +uniform float uSceneEpsilon; + +///////////////////////////////////////////////////////////////////////////////////////// +// Specific data types + +//! Stores ray parameters. +struct SRay +{ + vec3 Origin; + + vec3 Direct; +}; + +//! Stores intersection parameters. +struct SIntersect +{ + float Time; + + vec2 UV; + + vec3 Normal; +}; + +///////////////////////////////////////////////////////////////////////////////////////// +// Some useful constants + +#define MAXFLOAT 1e15f + +#define SMALL vec3 (exp2 (-80.f)) + +#define ZERO vec3 (0.f, 0.f, 0.f) +#define UNIT vec3 (1.f, 1.f, 1.f) + +#define AXIS_X vec3 (1.f, 0.f, 0.f) +#define AXIS_Y vec3 (0.f, 1.f, 0.f) +#define AXIS_Z vec3 (0.f, 0.f, 1.f) + +///////////////////////////////////////////////////////////////////////////////////////// +// Functions for compute ray-object intersection + +// ======================================================================= +// function : GenerateRay +// purpose : +// ======================================================================= +SRay GenerateRay (in vec2 thePixel) +{ + vec3 aP0 = mix (uOriginLB, uOriginRB, thePixel.x); + vec3 aP1 = mix (uOriginLT, uOriginRT, thePixel.x); + + vec3 aD0 = mix (uDirectLB, uDirectRB, thePixel.x); + vec3 aD1 = mix (uDirectLT, uDirectRT, thePixel.x); + + return SRay (mix (aP0, aP1, thePixel.y), + mix (aD0, aD1, thePixel.y)); +} + +// ======================================================================= +// function : IntersectSphere +// purpose : Computes ray-sphere intersection +// ======================================================================= +float IntersectSphere (in SRay theRay, in float theRadius) +{ + float aDdotD = dot (theRay.Direct, theRay.Direct); + float aDdotO = dot (theRay.Direct, theRay.Origin); + float aOdotO = dot (theRay.Origin, theRay.Origin); + + float aD = aDdotO * aDdotO - aDdotD * (aOdotO - theRadius * theRadius); + + if (aD > 0.f) + { + float aTime = (sqrt (aD) - aDdotO) * (1.f / aDdotD); + + return aTime > 0.f ? aTime : MAXFLOAT; + } + + return MAXFLOAT; +} + +// ======================================================================= +// function : IntersectTriangle +// purpose : Computes ray-triangle intersection (branchless version) +// ======================================================================= +float IntersectTriangle (in SRay theRay, + in vec3 thePnt0, + in vec3 thePnt1, + in vec3 thePnt2, + out vec2 theUV, + out vec3 theNorm) +{ + vec3 aEdge0 = thePnt1 - thePnt0; + vec3 aEdge1 = thePnt0 - thePnt2; + + theNorm = cross (aEdge1, aEdge0); + + vec3 aEdge2 = (1.f / dot (theNorm, theRay.Direct)) * (thePnt0 - theRay.Origin); + + float aTime = dot (theNorm, aEdge2); + + vec3 theVec = cross (theRay.Direct, aEdge2); + + theUV.x = dot (theVec, aEdge1); + theUV.y = dot (theVec, aEdge0); + + return bool (int(aTime >= 0.f) & + int(theUV.x >= 0.f) & + int(theUV.y >= 0.f) & + int(theUV.x + theUV.y <= 1.f)) ? aTime : MAXFLOAT; +} + +//! Global stack shared between traversal functions. +int Stack[STACK_SIZE]; + +//! Identifies the absence of intersection. +#define INALID_HIT ivec4 (-1) + +// ======================================================================= +// function : ObjectNearestHit +// purpose : Finds intersection with nearest object triangle +// ======================================================================= +ivec4 ObjectNearestHit (in int theBVHOffset, in int theVrtOffset, in int theTrgOffset, + in SRay theRay, in vec3 theInverse, inout SIntersect theHit, in int theSentinel) +{ + int aHead = theSentinel; // stack pointer + int aNode = 0; // node to visit + + ivec4 aTriIndex = INALID_HIT; + + float aTimeOut; + float aTimeLft; + float aTimeRgh; + + while (true) + { + ivec3 aData = texelFetch (uObjectNodeInfoTexture, aNode + theBVHOffset).xyz; + + if (aData.x == 0) // if inner node + { + vec3 aNodeMinLft = texelFetch (uObjectMinPointTexture, aData.y + theBVHOffset).xyz; + vec3 aNodeMaxLft = texelFetch (uObjectMaxPointTexture, aData.y + theBVHOffset).xyz; + vec3 aNodeMinRgh = texelFetch (uObjectMinPointTexture, aData.z + theBVHOffset).xyz; + vec3 aNodeMaxRgh = texelFetch (uObjectMaxPointTexture, aData.z + theBVHOffset).xyz; + + vec3 aTime0 = (aNodeMinLft - theRay.Origin) * theInverse; + vec3 aTime1 = (aNodeMaxLft - theRay.Origin) * theInverse; + + vec3 aTimeMax = max (aTime0, aTime1); + vec3 aTimeMin = min (aTime0, aTime1); + + aTime0 = (aNodeMinRgh - theRay.Origin) * theInverse; + aTime1 = (aNodeMaxRgh - theRay.Origin) * theInverse; + + aTimeOut = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z)); + aTimeLft = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z)); + + int aHitLft = int(aTimeLft <= aTimeOut) & int(aTimeOut >= 0.f) & int(aTimeLft <= theHit.Time); + + aTimeMax = max (aTime0, aTime1); + aTimeMin = min (aTime0, aTime1); + + aTimeOut = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z)); + aTimeRgh = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z)); + + int aHitRgh = int(aTimeRgh <= aTimeOut) & int(aTimeOut >= 0.f) & int(aTimeRgh <= theHit.Time); + + if (bool(aHitLft & aHitRgh)) + { + aNode = (aTimeLft < aTimeRgh) ? aData.y : aData.z; + + Stack[++aHead] = (aTimeLft < aTimeRgh) ? aData.z : aData.y; + } + else + { + if (bool(aHitLft | aHitRgh)) + { + aNode = bool(aHitLft) ? aData.y : aData.z; + } + else + { + if (aHead == theSentinel) + return aTriIndex; + + aNode = Stack[aHead--]; + } + } + } + else // if leaf node + { + vec3 aNormal; + vec2 aParams; + + for (int anIdx = aData.y; anIdx <= aData.z; ++anIdx) + { + ivec4 aTriangle = texelFetch (uGeometryTriangTexture, anIdx + theTrgOffset); + + vec3 aPoint0 = texelFetch (uGeometryVertexTexture, aTriangle.x + theVrtOffset).xyz; + vec3 aPoint1 = texelFetch (uGeometryVertexTexture, aTriangle.y + theVrtOffset).xyz; + vec3 aPoint2 = texelFetch (uGeometryVertexTexture, aTriangle.z + theVrtOffset).xyz; + + float aTime = IntersectTriangle (theRay, + aPoint0, + aPoint1, + aPoint2, + aParams, + aNormal); + + if (aTime < theHit.Time) + { + aTriIndex = aTriangle; + + theHit = SIntersect (aTime, aParams, aNormal); + } + } + + if (aHead == theSentinel) + return aTriIndex; + + aNode = Stack[aHead--]; + } + } + + return aTriIndex; +} + +// ======================================================================= +// function : ObjectAnyHit +// purpose : Finds intersection with any object triangle +// ======================================================================= +float ObjectAnyHit (in int theBVHOffset, in int theVrtOffset, in int theTrgOffset, + in SRay theRay, in vec3 theInverse, in float theDistance, in int theSentinel) +{ + int aHead = theSentinel; // stack pointer + int aNode = 0; // node to visit + + float aTimeOut; + float aTimeLft; + float aTimeRgh; + + while (true) + { + ivec4 aData = texelFetch (uObjectNodeInfoTexture, aNode + theBVHOffset); + + if (aData.x == 0) // if inner node + { + vec3 aNodeMinLft = texelFetch (uObjectMinPointTexture, aData.y + theBVHOffset).xyz; + vec3 aNodeMaxLft = texelFetch (uObjectMaxPointTexture, aData.y + theBVHOffset).xyz; + vec3 aNodeMinRgh = texelFetch (uObjectMinPointTexture, aData.z + theBVHOffset).xyz; + vec3 aNodeMaxRgh = texelFetch (uObjectMaxPointTexture, aData.z + theBVHOffset).xyz; + + vec3 aTime0 = (aNodeMinLft - theRay.Origin) * theInverse; + vec3 aTime1 = (aNodeMaxLft - theRay.Origin) * theInverse; + + vec3 aTimeMax = max (aTime0, aTime1); + vec3 aTimeMin = min (aTime0, aTime1); + + aTime0 = (aNodeMinRgh - theRay.Origin) * theInverse; + aTime1 = (aNodeMaxRgh - theRay.Origin) * theInverse; + + aTimeOut = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z)); + aTimeLft = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z)); + + int aHitLft = int(aTimeLft <= aTimeOut) & int(aTimeOut >= 0.f) & int(aTimeLft <= theDistance); + + aTimeMax = max (aTime0, aTime1); + aTimeMin = min (aTime0, aTime1); + + aTimeOut = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z)); + aTimeRgh = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z)); + + int aHitRgh = int(aTimeRgh <= aTimeOut) & int(aTimeOut >= 0.f) & int(aTimeRgh <= theDistance); + + if (bool(aHitLft & aHitRgh)) + { + aNode = (aTimeLft < aTimeRgh) ? aData.y : aData.z; + + Stack[++aHead] = (aTimeLft < aTimeRgh) ? aData.z : aData.y; + } + else + { + if (bool(aHitLft | aHitRgh)) + { + aNode = bool(aHitLft) ? aData.y : aData.z; + } + else + { + if (aHead == theSentinel) + return 1.f; + + aNode = Stack[aHead--]; + } + } + } + else // if leaf node + { + vec3 aNormal; + vec2 aParams; + + for (int anIdx = aData.y; anIdx <= aData.z; ++anIdx) + { + ivec4 aTriangle = texelFetch (uGeometryTriangTexture, anIdx + theTrgOffset); + + vec3 aPoint0 = texelFetch (uGeometryVertexTexture, aTriangle.x + theVrtOffset).xyz; + vec3 aPoint1 = texelFetch (uGeometryVertexTexture, aTriangle.y + theVrtOffset).xyz; + vec3 aPoint2 = texelFetch (uGeometryVertexTexture, aTriangle.z + theVrtOffset).xyz; + + float aTime = IntersectTriangle (theRay, + aPoint0, + aPoint1, + aPoint2, + aParams, + aNormal); + + if (aTime < theDistance) + return 0.f; + } + + if (aHead == theSentinel) + return 1.f; + + aNode = Stack[aHead--]; + } + } + + return 1.f; +} + +// ======================================================================= +// function : SceneNearestHit +// purpose : Finds intersection with nearest scene triangle +// ======================================================================= +ivec4 SceneNearestHit (in SRay theRay, in vec3 theInverse, inout SIntersect theHit) +{ + int aHead = -1; // stack pointer + int aNode = 0; // node to visit + + ivec4 aHitObject = INALID_HIT; + + float aTimeOut; + float aTimeLft; + float aTimeRgh; + + while (true) + { + ivec4 aData = texelFetch (uSceneNodeInfoTexture, aNode); + + if (aData.x != 0) // if leaf node + { + vec3 aNodeMin = texelFetch (uSceneMinPointTexture, aNode).xyz; + vec3 aNodeMax = texelFetch (uSceneMaxPointTexture, aNode).xyz; + + vec3 aTime0 = (aNodeMin - theRay.Origin) * theInverse; + vec3 aTime1 = (aNodeMax - theRay.Origin) * theInverse; + + vec3 aTimes = min (aTime0, aTime1); + + if (max (aTimes.x, max (aTimes.y, aTimes.z)) < theHit.Time) + { + ivec4 aTriIndex = ObjectNearestHit ( + aData.y, aData.z, aData.w, theRay, theInverse, theHit, aHead); + + if (aTriIndex.x != -1) + { + aHitObject = ivec4 (aTriIndex.x + aData.z, // vertex 0 + aTriIndex.y + aData.z, // vertex 1 + aTriIndex.z + aData.z, // vertex 2 + aTriIndex.w); // material + } + } + + if (aHead < 0) + return aHitObject; + + aNode = Stack[aHead--]; + } + else // if inner node + { + vec3 aNodeMinLft = texelFetch (uSceneMinPointTexture, aData.y).xyz; + vec3 aNodeMaxLft = texelFetch (uSceneMaxPointTexture, aData.y).xyz; + vec3 aNodeMinRgh = texelFetch (uSceneMinPointTexture, aData.z).xyz; + vec3 aNodeMaxRgh = texelFetch (uSceneMaxPointTexture, aData.z).xyz; + + vec3 aTime0 = (aNodeMinLft - theRay.Origin) * theInverse; + vec3 aTime1 = (aNodeMaxLft - theRay.Origin) * theInverse; + + vec3 aTimeMax = max (aTime0, aTime1); + vec3 aTimeMin = min (aTime0, aTime1); + + aTimeOut = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z)); + aTimeLft = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z)); + + int aHitLft = int(aTimeLft <= aTimeOut) & int(aTimeOut >= 0.f) & int(aTimeLft <= theHit.Time); + + aTime0 = (aNodeMinRgh - theRay.Origin) * theInverse; + aTime1 = (aNodeMaxRgh - theRay.Origin) * theInverse; + + aTimeMax = max (aTime0, aTime1); + aTimeMin = min (aTime0, aTime1); + + aTimeOut = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z)); + aTimeRgh = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z)); + + int aHitRgh = int(aTimeRgh <= aTimeOut) & int(aTimeOut >= 0.f) & int(aTimeRgh <= theHit.Time); + + if (bool(aHitLft & aHitRgh)) + { + aNode = (aTimeLft < aTimeRgh) ? aData.y : aData.z; + + Stack[++aHead] = (aTimeLft < aTimeRgh) ? aData.z : aData.y; + } + else + { + if (bool(aHitLft | aHitRgh)) + { + aNode = bool(aHitLft) ? aData.y : aData.z; + } + else + { + if (aHead < 0) + return aHitObject; + + aNode = Stack[aHead--]; + } + } + } + } + + return aHitObject; +} + +// ======================================================================= +// function : SceneAnyHit +// purpose : Finds intersection with any scene triangle +// ======================================================================= +float SceneAnyHit (in SRay theRay, in vec3 theInverse, in float theDistance) +{ + int aHead = -1; // stack pointer + int aNode = 0; // node to visit + + float aTimeOut; + float aTimeLft; + float aTimeRgh; + + while (true) + { + ivec4 aData = texelFetch (uSceneNodeInfoTexture, aNode); + + if (aData.x != 0) // if leaf node + { + bool isShadow = 0.f == ObjectAnyHit ( + aData.y, aData.z, aData.w, theRay, theInverse, theDistance, aHead); + + if (aHead < 0 || isShadow) + return isShadow ? 0.f : 1.f; + + aNode = Stack[aHead--]; + } + else // if inner node + { + vec3 aNodeMinLft = texelFetch (uSceneMinPointTexture, aData.y).xyz; + vec3 aNodeMaxLft = texelFetch (uSceneMaxPointTexture, aData.y).xyz; + vec3 aNodeMinRgh = texelFetch (uSceneMinPointTexture, aData.z).xyz; + vec3 aNodeMaxRgh = texelFetch (uSceneMaxPointTexture, aData.z).xyz; + + vec3 aTime0 = (aNodeMinLft - theRay.Origin) * theInverse; + vec3 aTime1 = (aNodeMaxLft - theRay.Origin) * theInverse; + + vec3 aTimeMax = max (aTime0, aTime1); + vec3 aTimeMin = min (aTime0, aTime1); + + aTimeOut = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z)); + aTimeLft = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z)); + + int aHitLft = int(aTimeLft <= aTimeOut) & int(aTimeOut >= 0.f) & int(aTimeLft <= theDistance); + + aTime0 = (aNodeMinRgh - theRay.Origin) * theInverse; + aTime1 = (aNodeMaxRgh - theRay.Origin) * theInverse; + + aTimeMax = max (aTime0, aTime1); + aTimeMin = min (aTime0, aTime1); + + aTimeOut = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z)); + aTimeRgh = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z)); + + int aHitRgh = int(aTimeRgh <= aTimeOut) & int(aTimeOut >= 0.f) & int(aTimeRgh <= theDistance); + + if (bool(aHitLft & aHitRgh)) + { + aNode = (aTimeLft < aTimeRgh) ? aData.y : aData.z; + + Stack[++aHead] = (aTimeLft < aTimeRgh) ? aData.z : aData.y; + } + else + { + if (bool(aHitLft | aHitRgh)) + { + aNode = bool(aHitLft) ? aData.y : aData.z; + } + else + { + if (aHead < 0) + return 1.f; + + aNode = Stack[aHead--]; + } + } + } + } + + return 1.f; +} + +#define PI 3.1415926f + +// ======================================================================= +// function : Latlong +// purpose : Converts world direction to environment texture coordinates +// ======================================================================= +vec2 Latlong (in vec3 thePoint, in float theRadius) +{ + float aPsi = acos (-thePoint.z / theRadius); + + float aPhi = atan (thePoint.y, thePoint.x) + PI; + + return vec2 (aPhi * 0.1591549f, + aPsi * 0.3183098f); +} + +// ======================================================================= +// function : SmoothNormal +// purpose : Interpolates normal across the triangle +// ======================================================================= +vec3 SmoothNormal (in vec2 theUV, in ivec4 theTriangle) +{ + vec3 aNormal0 = texelFetch (uGeometryNormalTexture, theTriangle.x).xyz; + vec3 aNormal1 = texelFetch (uGeometryNormalTexture, theTriangle.y).xyz; + vec3 aNormal2 = texelFetch (uGeometryNormalTexture, theTriangle.z).xyz; + + return normalize (aNormal1 * theUV.x + + aNormal2 * theUV.y + + aNormal0 * (1.f - theUV.x - theUV.y)); +} + +#define THRESHOLD vec3 (0.1f, 0.1f, 0.1f) + +#define MATERIAL_AMBN(index) (7 * index + 0) +#define MATERIAL_DIFF(index) (7 * index + 1) +#define MATERIAL_SPEC(index) (7 * index + 2) +#define MATERIAL_EMIS(index) (7 * index + 3) +#define MATERIAL_REFL(index) (7 * index + 4) +#define MATERIAL_REFR(index) (7 * index + 5) +#define MATERIAL_TRAN(index) (7 * index + 6) + +#define LIGHT_POS(index) (2 * index + 1) +#define LIGHT_PWR(index) (2 * index + 0) + +// ======================================================================= +// function : Radiance +// purpose : Computes color of specified ray +// ======================================================================= +vec4 Radiance (in SRay theRay, in vec3 theInverse) +{ + vec3 aResult = vec3 (0.f); + vec4 aWeight = vec4 (1.f); + + for (int aDepth = 0; aDepth < 5; ++aDepth) + { + SIntersect aHit = SIntersect (MAXFLOAT, vec2 (ZERO), ZERO); + + ivec4 aTriIndex = SceneNearestHit (theRay, theInverse, aHit); + + if (aTriIndex.x == -1) + { + if (aWeight.w != 0.f) + { + return vec4 (aResult.x, + aResult.y, + aResult.z, + aWeight.w); + } + + if (bool(uEnvironmentEnable)) + { + float aTime = IntersectSphere (theRay, uSceneRadius); + + aResult.xyz += aWeight.xyz * textureLod (uEnvironmentMapTexture, + Latlong (theRay.Direct * aTime + theRay.Origin, uSceneRadius), 0.f).xyz; + } + + return vec4 (aResult.x, + aResult.y, + aResult.z, + aWeight.w); + } + + vec3 aPoint = theRay.Direct * aHit.Time + theRay.Origin; + + vec3 aAmbient = vec3 (texelFetch ( + uRaytraceMaterialTexture, MATERIAL_AMBN (aTriIndex.w))); + vec3 aDiffuse = vec3 (texelFetch ( + uRaytraceMaterialTexture, MATERIAL_DIFF (aTriIndex.w))); + vec4 aSpecular = vec4 (texelFetch ( + uRaytraceMaterialTexture, MATERIAL_SPEC (aTriIndex.w))); + vec2 aOpacity = vec2 (texelFetch ( + uRaytraceMaterialTexture, MATERIAL_TRAN (aTriIndex.w))); + + vec3 aNormal = SmoothNormal (aHit.UV, aTriIndex); + + aHit.Normal = normalize (aHit.Normal); + + for (int aLightIdx = 0; aLightIdx < uLightCount; ++aLightIdx) + { + vec4 aLight = texelFetch ( + uRaytraceLightSrcTexture, LIGHT_POS (aLightIdx)); + + float aDistance = MAXFLOAT; + + if (aLight.w != 0.f) // point light source + { + aDistance = length (aLight.xyz -= aPoint); + + aLight.xyz *= 1.f / aDistance; + } + + SRay aShadow = SRay (aPoint + aLight.xyz * uSceneEpsilon, aLight.xyz); + + aShadow.Origin += aHit.Normal * uSceneEpsilon * + (dot (aHit.Normal, aLight.xyz) >= 0.f ? 1.f : -1.f); + + float aVisibility = 1.f; + + if (bool(uShadowsEnable)) + { + vec3 aInverse = 1.f / max (abs (aLight.xyz), SMALL); + + aInverse.x = aLight.x < 0.f ? -aInverse.x : aInverse.x; + aInverse.y = aLight.y < 0.f ? -aInverse.y : aInverse.y; + aInverse.z = aLight.z < 0.f ? -aInverse.z : aInverse.z; + + aVisibility = SceneAnyHit (aShadow, aInverse, aDistance); + } + + if (aVisibility > 0.f) + { + vec3 aIntensity = vec3 (texelFetch ( + uRaytraceLightSrcTexture, LIGHT_PWR (aLightIdx))); + + float aLdotN = dot (aShadow.Direct, aNormal); + + if (aOpacity.y > 0.f) // force two-sided lighting + aLdotN = abs (aLdotN); // for transparent surfaces + + if (aLdotN > 0.f) + { + float aRdotV = dot (reflect (aShadow.Direct, aNormal), theRay.Direct); + + aResult.xyz += aWeight.xyz * aOpacity.x * aIntensity * + (aDiffuse * aLdotN + aSpecular.xyz * pow (max (0.f, aRdotV), aSpecular.w)); + } + } + } + + aResult.xyz += aWeight.xyz * uGlobalAmbient.xyz * + aAmbient * aOpacity.x * max (abs (dot (aNormal, theRay.Direct)), 0.5f); + + if (aOpacity.x != 1.f) + { + aWeight *= aOpacity.y; + } + else + { + aWeight *= bool(uReflectionsEnable) ? + texelFetch (uRaytraceMaterialTexture, MATERIAL_REFL (aTriIndex.w)) : vec4 (0.f); + + theRay.Direct = reflect (theRay.Direct, aNormal); + + if (dot (theRay.Direct, aHit.Normal) < 0.f) + { + theRay.Direct = reflect (theRay.Direct, aHit.Normal); + } + + theInverse = 1.0 / max (abs (theRay.Direct), SMALL); + + theInverse.x = theRay.Direct.x < 0.0 ? -theInverse.x : theInverse.x; + theInverse.y = theRay.Direct.y < 0.0 ? -theInverse.y : theInverse.y; + theInverse.z = theRay.Direct.z < 0.0 ? -theInverse.z : theInverse.z; + + aPoint += aHit.Normal * (dot (aHit.Normal, theRay.Direct) >= 0.f ? uSceneEpsilon : -uSceneEpsilon); + } + + if (all (lessThanEqual (aWeight.xyz, THRESHOLD))) + { + return vec4 (aResult.x, + aResult.y, + aResult.z, + aWeight.w); + } + + theRay.Origin = theRay.Direct * uSceneEpsilon + aPoint; + } + + return vec4 (aResult.x, + aResult.y, + aResult.z, + aWeight.w); +} diff --git a/src/Shaders/RaytraceBase.vs b/src/Shaders/RaytraceBase.vs new file mode 100644 index 0000000000..6ff4d01a6c --- /dev/null +++ b/src/Shaders/RaytraceBase.vs @@ -0,0 +1,12 @@ +in vec4 aPosition; + +//! Normalized pixel coordinates. +out vec2 vPixel; + +void main (void) +{ + vPixel = vec2 ((aPosition.x + 1.f) * 0.5f, + (aPosition.y + 1.f) * 0.5f); + + gl_Position = aPosition; +} \ No newline at end of file diff --git a/src/Shaders/RaytraceRender.fs b/src/Shaders/RaytraceRender.fs new file mode 100644 index 0000000000..76100664cd --- /dev/null +++ b/src/Shaders/RaytraceRender.fs @@ -0,0 +1,18 @@ +out vec4 OutColor; + +// ======================================================================= +// function : main +// purpose : +// ======================================================================= +void main (void) +{ + SRay aRay = GenerateRay (vPixel); + + vec3 aInvDirect = 1.f / max (abs (aRay.Direct), SMALL); + + aInvDirect = vec3 (aRay.Direct.x < 0.f ? -aInvDirect.x : aInvDirect.x, + aRay.Direct.y < 0.f ? -aInvDirect.y : aInvDirect.y, + aRay.Direct.z < 0.f ? -aInvDirect.z : aInvDirect.z); + + OutColor = clamp (Radiance (aRay, aInvDirect), 0.f, 1.f); +} \ No newline at end of file diff --git a/src/Shaders/RaytraceSmooth.fs b/src/Shaders/RaytraceSmooth.fs new file mode 100644 index 0000000000..f75a971799 --- /dev/null +++ b/src/Shaders/RaytraceSmooth.fs @@ -0,0 +1,79 @@ +//! Input ray-traced image. +uniform sampler2D uFSAAInputTexture; + +//! Number of accumulated FSAA samples. +uniform int uSamples; + +//! Sub-pixel offset in X direction for FSAA. +uniform float uOffsetX; +//! Sub-pixel offset in Y direction for FSAA. +uniform float uOffsetY; + +//! Output pixel color. +out vec4 OutColor; + +#define LUM_DIFFERENCE 0.085f + +#define LUMA vec3 (0.2126f, 0.7152f, 0.0722f) + +// ======================================================================= +// function : main +// purpose : +// ======================================================================= +void main (void) +{ + int aPixelX = int (gl_FragCoord.x); + int aPixelY = int (gl_FragCoord.y); + + vec4 aClr0 = texelFetch (uFSAAInputTexture, ivec2 (aPixelX + 0, aPixelY + 0), 0); + vec4 aClr1 = texelFetch (uFSAAInputTexture, ivec2 (aPixelX + 0, aPixelY - 1), 0); + vec4 aClr2 = texelFetch (uFSAAInputTexture, ivec2 (aPixelX + 0, aPixelY + 1), 0); + + vec4 aClr3 = texelFetch (uFSAAInputTexture, ivec2 (aPixelX + 1, aPixelY + 0), 0); + vec4 aClr4 = texelFetch (uFSAAInputTexture, ivec2 (aPixelX + 1, aPixelY - 1), 0); + vec4 aClr5 = texelFetch (uFSAAInputTexture, ivec2 (aPixelX + 1, aPixelY + 1), 0); + + vec4 aClr6 = texelFetch (uFSAAInputTexture, ivec2 (aPixelX - 1, aPixelY + 0), 0); + vec4 aClr7 = texelFetch (uFSAAInputTexture, ivec2 (aPixelX - 1, aPixelY - 1), 0); + vec4 aClr8 = texelFetch (uFSAAInputTexture, ivec2 (aPixelX - 1, aPixelY + 1), 0); + + float aLum = dot (LUMA, aClr0.xyz); + + bool aRender = abs (aClr1.w - aClr0.w) > LUM_DIFFERENCE || + abs (aClr2.w - aClr0.w) > LUM_DIFFERENCE || + abs (aClr3.w - aClr0.w) > LUM_DIFFERENCE || + abs (aClr4.w - aClr0.w) > LUM_DIFFERENCE || + abs (aClr5.w - aClr0.w) > LUM_DIFFERENCE || + abs (aClr6.w - aClr0.w) > LUM_DIFFERENCE || + abs (aClr7.w - aClr0.w) > LUM_DIFFERENCE || + abs (aClr8.w - aClr0.w) > LUM_DIFFERENCE; + + if (!aRender) + { + aRender = abs (dot (LUMA, aClr1.xyz) - aLum) > LUM_DIFFERENCE || + abs (dot (LUMA, aClr2.xyz) - aLum) > LUM_DIFFERENCE || + abs (dot (LUMA, aClr3.xyz) - aLum) > LUM_DIFFERENCE || + abs (dot (LUMA, aClr4.xyz) - aLum) > LUM_DIFFERENCE || + abs (dot (LUMA, aClr5.xyz) - aLum) > LUM_DIFFERENCE || + abs (dot (LUMA, aClr6.xyz) - aLum) > LUM_DIFFERENCE || + abs (dot (LUMA, aClr7.xyz) - aLum) > LUM_DIFFERENCE || + abs (dot (LUMA, aClr8.xyz) - aLum) > LUM_DIFFERENCE; + } + + vec4 aColor = aClr0; + + if (aRender) + { + SRay aRay = GenerateRay (vPixel + vec2 (uOffsetX, uOffsetY)); + + vec3 aInvDirect = 1.f / max (abs (aRay.Direct), SMALL); + + aInvDirect = vec3 (aRay.Direct.x < 0.f ? -aInvDirect.x : aInvDirect.x, + aRay.Direct.y < 0.f ? -aInvDirect.y : aInvDirect.y, + aRay.Direct.z < 0.f ? -aInvDirect.z : aInvDirect.z); + + aColor = mix (aClr0, clamp (Radiance (aRay, aInvDirect), 0.f, 1.f), 1.f / uSamples); + } + + OutColor = aColor; +} \ No newline at end of file diff --git a/src/TKOpenGl/EXTERNLIB b/src/TKOpenGl/EXTERNLIB index a27a27af32..078b54bbbf 100755 --- a/src/TKOpenGl/EXTERNLIB +++ b/src/TKOpenGl/EXTERNLIB @@ -6,7 +6,6 @@ CSF_OpenGlLibs CSF_objc CSF_Appkit CSF_IOKit -CSF_OPENCL CSF_FREETYPE CSF_GL2PS CSF_user32 diff --git a/src/ViewerTest/ViewerTest_ViewerCommands.cxx b/src/ViewerTest/ViewerTest_ViewerCommands.cxx index f4c5cfec2e..b002336435 100644 --- a/src/ViewerTest/ViewerTest_ViewerCommands.cxx +++ b/src/ViewerTest/ViewerTest_ViewerCommands.cxx @@ -6288,44 +6288,6 @@ static int VLight (Draw_Interpretor& theDi, return 0; } - - -//============================================================================== -//function : VClInfo -//purpose : Prints info about active OpenCL device -//============================================================================== - -static Standard_Integer VClInfo (Draw_Interpretor& theDi, - Standard_Integer, - const char**) -{ - Handle(AIS_InteractiveContext) aContextAIS = ViewerTest::GetAISContext(); - if (aContextAIS.IsNull()) - { - std::cerr << "Call vinit before!\n"; - return 1; - } - - Handle(OpenGl_GraphicDriver) aDrv = Handle(OpenGl_GraphicDriver)::DownCast (aContextAIS->CurrentViewer()->Driver()); - Graphic3d_CView* aCView = static_cast (ViewerTest::CurrentView()->View()->CView()); - NCollection_DataMap anInfo; - if (aDrv.IsNull() - || aCView == NULL - || !aDrv->GetOpenClDeviceInfo (*aCView, anInfo)) - { - theDi << "OpenCL device info is unavailable!\n"; - return 0; - } - - theDi << "OpenCL device info:\n"; - for (NCollection_DataMap::Iterator anIter (anInfo); - anIter.More(); anIter.Next()) - { - theDi << anIter.Key() << ": \t" << anIter.Value() << "\n"; - } - return 0; -} - //======================================================================= //function : VRaytrace //purpose : Enables/disables OpenCL-based ray-tracing @@ -6756,9 +6718,6 @@ void ViewerTest::ViewerCommands(Draw_Interpretor& theCommands) theCommands.Add("vraytrace", "vraytrace 0|1", __FILE__,VRaytrace,group); - theCommands.Add("vclinfo", - "vclinfo", - __FILE__,VClInfo,group); theCommands.Add("vsetraytracemode", "vsetraytracemode [shad=0|1] [refl=0|1] [aa=0|1]", __FILE__,VSetRaytraceMode,group); diff --git a/tests/v3d/raytrace/bug24130 b/tests/v3d/raytrace/bug24130 index c3b41b63d2..2837eec3b2 100644 --- a/tests/v3d/raytrace/bug24130 +++ b/tests/v3d/raytrace/bug24130 @@ -23,7 +23,6 @@ vfit # activate ray-tracing vraytrace 1 -vclinfo set aModeNum 0 for { set aAAMode 0 } { $aAAMode <= 1 } { incr aAAMode } { diff --git a/tests/v3d/raytrace/connected b/tests/v3d/raytrace/connected index 4ea2525f7d..3d67b6076b 100644 --- a/tests/v3d/raytrace/connected +++ b/tests/v3d/raytrace/connected @@ -24,7 +24,6 @@ vdump $::imagedir/${::casename}_OFF.png # turn on ray tracing vraytrace 1 -vclinfo vdump $::imagedir/${::casename}_rt1.png vclear diff --git a/tests/v3d/raytrace/plastic b/tests/v3d/raytrace/plastic index 54c08d004b..b46b862c35 100644 --- a/tests/v3d/raytrace/plastic +++ b/tests/v3d/raytrace/plastic @@ -25,7 +25,6 @@ if { "$aColorL" != "GREEN3" || "$aColorR" != "GREEN4" } { } vraytrace 1 -vclinfo set aColorL [vreadpixel 150 250 rgb name] set aColorR [vreadpixel 250 250 rgb name] #if { "$aColorL" != "GREEN3" || "$aColorR" != "GREEN4" } {