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
}
|| theBVH->Level (aChildIndex) >= BVH_Builder<T, N>::myMaxTreeDepth;
if (!isLeaf)
+ {
BVH_Builder<T, N>::myTasksQueue.Append (aChildIndex);
+ }
+
+ BVH_Builder<T, N>::UpdateDepth (theBVH, theBVH->Level (aChildIndex));
}
}
BVH_Tree<T, N>* theBVH,
const Standard_Integer theTask);
+ //! Updates depth of constructed BVH tree.
+ void UpdateDepth (BVH_Tree<T, N>* theBVH,
+ const Standard_Integer theLevel)
+ {
+ if (theLevel > theBVH->myDepth)
+ {
+ theBVH->myDepth = theLevel;
+ }
+ }
+
protected:
Standard_Integer myMaxTreeDepth; //!< Maximum depth of constructed BVH
{
BVH_Builder<T, N>::myTasksQueue.Append (aChildIndex);
}
+
+ BVH_Builder<T, N>::UpdateDepth (theBVH, theBVH->Level (aChildIndex));
}
}
#include <BVH_Box.hxx>
+template<class T, int N> 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
template<class T, int N>
class BVH_Tree
{
+ friend class BVH_Builder<T, N>;
+
public:
typedef typename BVH_Box<T, N>::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)
{
return BVHTools::ArrayOp<Standard_Integer, 4>::Size (myNodeInfoBuffer);
}
+ //! Returns depth of BVH tree from last build.
+ Standard_Integer Depth() const
+ {
+ return myDepth;
+ }
+
public:
//! Removes all BVH nodes.
//! Array of node data records.
BVH_Array4i myNodeInfoBuffer;
+ //! Depth of constructed tree.
+ Standard_Integer myDepth;
+
};
#include <BVH_Tree.lxx>
template<class T, int N>
void BVH_Tree<T, N>::Clear()
{
+ myDepth = 0;
+
BVHTools::ArrayOp<T, N>::Clear (myMinPointBuffer);
BVHTools::ArrayOp<T, N>::Clear (myMaxPointBuffer);
WasRedrawnGL (0),
IsRaytracing (0),
IsShadowsEnabled (1),
- IsReflectionsEnabled (1),
+ IsReflectionsEnabled (0),
IsAntialiasingEnabled (0)
{
memset(&DefWindow,0,sizeof(DefWindow));
CSF_Appkit
CSF_IOKit
CSF_OpenGlLibs
-CSF_OPENCL
CSF_AviLibs
CSF_FREETYPE
CSF_GL2PS
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
+++ /dev/null
-// 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 <OpenGl_GlCore20.hxx>
-
-#if defined(__APPLE__) || defined(__MACOSX)
- #include <OpenCL/opencl.h>
-#else
- #include <CL/cl.h>
- #include <CL/cl_gl.h>
-#endif
-
-#endif // _OpenGl_Cl_H__
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<TCollection_AsciiString, TCollection_AsciiString>&)
-{
- return Standard_False;
-}
-
-#else
-
-Standard_Boolean OpenGl_GraphicDriver::GetOpenClDeviceInfo (const Graphic3d_CView& theCView,
- NCollection_DataMap<TCollection_AsciiString, TCollection_AsciiString>& theInfo)
-{
-
- if (theCView.ViewId == -1 || theCView.ptrView == NULL)
- {
- return Standard_False;
- }
-
- return reinterpret_cast<const OpenGl_CView*> (theCView.ptrView)->WS->GetOpenClDeviceInfo (theInfo);
-}
-
-#endif
-
// =======================================================================
// function : DisplayImmediateStructure
// purpose :
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<TCollection_AsciiString, TCollection_AsciiString>& theInfo);
//! Method to retrieve valid GL context.
//! Could return NULL-handle if no window created by this driver.
// Alternatively, this file may be used under the terms of Open CASCADE
// commercial license or contractual agreement.
-#ifdef HAVE_CONFIG_H
- #include <config.h>
-#endif
-
#include <OpenGl_Group.hxx>
#include <OpenGl_GraphicDriver.hxx>
AddElement (anAspectFace);
}
-#ifdef HAVE_OPENCL
if (myIsRaytracable)
{
++myModificationState;
aStruct->UpdateStateWithAncestorStructures();
}
}
-#endif
}
// =======================================================================
(myLast? myLast->next : myFirst) = aNode;
myLast = aNode;
-#ifdef HAVE_OPENCL
if (OpenGl_Raytrace::IsRaytracedElement (aNode))
{
myModificationState++;
aStruct->SetRaytracableWithAncestorStructures();
}
}
-#endif
}
// =======================================================================
// Alternatively, this file may be used under the terms of Open CASCADE
// commercial license or contractual agreement.
-#ifdef HAVE_CONFIG_H
- #include <config.h>
-#endif
-
#include <OpenGl_GlCore11.hxx>
#include <OpenGl_LayerList.hxx>
{
myNbStructures--;
-#ifdef HAVE_OPENCL
if (theStructure->IsRaytracable())
{
myModificationState++;
}
-#endif
return;
}
{
myNbStructures--;
-#ifdef HAVE_OPENCL
if (theStructure->IsRaytracable())
{
myModificationState++;
}
-#endif
return;
}
//! 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
Standard_Integer myNbPriorities;
Standard_Integer myNbStructures;
-#ifdef HAVE_OPENCL
mutable Standard_Size myModificationState;
-#endif
public:
DEFINE_STANDARD_ALLOC
+++ /dev/null
-// 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 <config.h>
-#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
// Alternatively, this file may be used under the terms of Open CASCADE
// commercial license or contractual agreement.
-#ifdef HAVE_CONFIG_H
- #include <config.h>
-#endif
-
-#ifdef HAVE_OPENCL
-
#include <Standard_Assert.hxx>
#ifdef HAVE_TBB
OpenGL_BVHParallelBuilder (this));
#endif
+ myBottomLevelTreeDepth = 0;
+
for (Standard_Integer anObjectIdx = 0; anObjectIdx < Size(); ++anObjectIdx)
{
OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
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
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;
return Standard_False;
}
}
-
-#endif
#ifndef _OpenGl_SceneGeometry_Header
#define _OpenGl_SceneGeometry_Header
-#ifdef HAVE_OPENCL
-
#include <BVH_Geometry.hxx>
#include <BVH_Triangulation.hxx>
#include <NCollection_StdAllocator.hxx>
{
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<Standard_ShortReal, 4>()
{
//
}
NCollection_StdAllocator<OpenGl_RaytraceMaterial> > Materials;
//! Global ambient from all light sources.
- BVH_Vec4f GlobalAmbient;
+ BVH_Vec4f Ambient;
public:
//! Creates uninitialized ray-tracing geometry.
OpenGl_RaytraceGeometry()
+ : BVH_Geometry<Standard_ShortReal, 4>(),
+ myHighLevelTreeDepth (0),
+ myBottomLevelTreeDepth (0)
{
//
}
//! 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
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;
}
// =======================================================================
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
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
//! Wrapper for OpenGL program object.
class OpenGl_ShaderProgram : public OpenGl_Resource
{
+ friend class OpenGl_Workspace;
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.
// Alternatively, this file may be used under the terms of Open CASCADE
// commercial license or contractual agreement.
-#ifdef HAVE_CONFIG_H
- #include <config.h>
-#endif
-
#include <OpenGl_CappingAlgo.hxx>
#include <OpenGl_Context.hxx>
#include <OpenGl_GlCore11.hxx>
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
}
// =======================================================================
matcpy (myTransformation->mat, &Graphic3d_CStructure::Transformation[0][0]);
-#ifdef HAVE_OPENCL
if (myIsRaytracable)
{
UpdateStateWithAncestorStructures();
}
-#endif
}
// =======================================================================
}
myAspectFace->SetAspect (theAspect);
-#ifdef HAVE_OPENCL
if (myIsRaytracable)
{
UpdateStateWithAncestorStructures();
}
-#endif
}
// =======================================================================
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 :
}
}
-#endif
-
// =======================================================================
// function : Connect
// purpose :
Disconnect (theStructure);
myConnected.Append (aStruct);
-#ifdef HAVE_OPENCL
if (aStruct->IsRaytracable())
{
UpdateStateWithAncestorStructures();
}
aStruct->RegisterAncestorStructure (this);
-#endif
}
// =======================================================================
{
myConnected.Remove (anIter);
-#ifdef HAVE_OPENCL
if (aStruct->IsRaytracable())
{
UpdateStateWithAncestorStructures();
}
aStruct->UnregisterAncestorStructure (this);
-#endif
return;
}
}
{
theGroup->Clear (Standard_False);
- #ifdef HAVE_OPENCL
if (((OpenGl_Group* )theGroup.operator->())->IsRaytracable())
{
UpdateStateWithAncestorStructures();
UpdateRaytracableWithAncestorStructures();
}
- #endif
myGroups.Remove (aGroupIter);
return;
// =======================================================================
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
}
// =======================================================================
OpenGl_Element::Destroy (theGlCtx, myAspectText);
clearHighlightColor (theGlCtx);
-#ifdef HAVE_OPENCL
// Remove from connected list of ancestor
UnregisterFromAncestorStructure();
-#endif
}
// =======================================================================
//! 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; }
//! 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;
//! Sets ray-tracable status for structure and its parents.
void SetRaytracableWithAncestorStructures() const;
-#endif
-
protected:
OpenGl_Matrix* myTransformation;
OpenGl_ListOfStructure myConnected;
-#ifdef HAVE_OPENCL
mutable OpenGl_ListOfStructure myAncestorStructures;
mutable Standard_Boolean myIsRaytracable;
mutable Standard_Size myModificationState;
-#endif
public:
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 :
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;
// Alternatively, this file may be used under the terms of Open CASCADE
// commercial license or contractual agreement.
-#ifdef HAVE_CONFIG_H
- #include <config.h>
-#endif
-
#include <NCollection_Mat4.hxx>
#include <OpenGl_Context.hxx>
myCurrLightSourceState = myStateCounter->Increment();
-#ifdef HAVE_OPENCL
myModificationState = 1; // initial state
-#endif
}
/*----------------------------------------------------------------------*/
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
}
// =======================================================================
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; }
return myImmediateList;
}
-#ifdef HAVE_OPENCL
//! Returns modification state for ray-tracing.
Standard_Size ModificationState() const { return myModificationState; }
-#endif
protected:
StateInfo myLastViewMappingState;
StateInfo myLastLightSourceState;
-#ifdef HAVE_OPENCL
Standard_Size myModificationState;
-#endif
public:
/*----------------------------------------------------------------------*/
-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 );
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
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);
// ====================================
// Render background
- DrawBackground (theWorkspace);
+ DrawBackground (*theWorkspace);
// Switch off lighting by default
glDisable(GL_LIGHTING);
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),
// Polygon Offset
EnablePolygonOffset();
-
-#ifdef HAVE_OPENCL
-
- myComputeInitStatus = OpenGl_CLIS_NONE;
-
- myViewModificationStatus = 0;
- myLayersModificationStatus = 0;
-
- myIsRaytraceDataValid = Standard_False;
- myToUpdateRaytraceData = Standard_False;
-
-#endif
}
// =======================================================================
// =======================================================================
OpenGl_Workspace::~OpenGl_Workspace()
{
-#ifdef HAVE_OPENCL
- ReleaseOpenCL();
-#endif
+ ReleaseRaytraceResources();
}
// =======================================================================
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)
}
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)
{
#ifndef _OpenGl_Workspace_Header
#define _OpenGl_Workspace_Header
-#ifdef HAVE_OPENCL
- #include <map>
- #include <set>
-
- #include <OpenGl_Cl.hxx>
-#endif
+#include <map>
+#include <set>
#include <Handle_OpenGl_Workspace.hxx>
#include <OpenGl_Window.hxx>
#include <OpenGl_AspectFace.hxx>
#include <OpenGl_Display.hxx>
+#include <OpenGl_FrameBuffer.hxx>
#include <OpenGl_Matrix.hxx>
#include <OpenGl_NamedStatus.hxx>
#include <OpenGl_PrinterContext.hxx>
-#ifdef HAVE_OPENCL
- #include <OpenGl_SceneGeometry.hxx>
-#endif
+#include <OpenGl_SceneGeometry.hxx>
#include <OpenGl_TextParam.hxx>
#include <OpenGl_RenderFilter.hxx>
#include <OpenGl_Vec.hxx>
#include <Handle_OpenGl_View.hxx>
#include <Handle_OpenGl_Texture.hxx>
+#include <OpenGl_ShaderObject.hxx>
+#include <OpenGl_ShaderProgram.hxx>
+#include <OpenGl_TextureBufferArb.hxx>
+
class OpenGl_AspectLine;
class OpenGl_AspectMarker;
class OpenGl_AspectText;
void setTextureParams (Handle(OpenGl_Texture)& theTexture,
const Handle(Graphic3d_TextureParams)& theParams);
-#ifdef HAVE_OPENCL
+protected:
-public:
+ //! Result of OpenGL shaders initialization.
+ enum RaytraceInitStatus
+ {
+ OpenGl_RT_NONE,
+ OpenGl_RT_INIT,
+ OpenGl_RT_FAIL
+ };
- //! Returns information about OpenCL device used for computations.
- Standard_Boolean GetOpenClDeviceInfo (
- NCollection_DataMap<TCollection_AsciiString, TCollection_AsciiString>& theInfo) const;
+ //! 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
+ };
-protected:
+ //! 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,
- //! Describes result of OpenCL initializing.
- enum OpenClInitStatus
+ 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
{
- OpenGl_CLIS_NONE,
- OpenGl_CLIS_INIT,
- OpenGl_CLIS_FAIL
+ 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.
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);
+
+ //! Performs safe exit when shaders initialization fails.
+ Standard_Boolean SafeFailBack (const TCollection_ExtendedString& theMessage);
+
+ //! Initializes OpenGL/GLSL shader programs.
+ Standard_Boolean InitRaytraceResources();
- //! Releases OpenCL resources.
- void ReleaseOpenCL();
+ //! Releases OpenGL/GLSL shader programs.
+ void ReleaseRaytraceResources();
- //! Resizes OpenCL output image.
- Standard_Boolean ResizeRaytraceOutputBuffer (const cl_int theSizeX, const cl_int theSizeY);
+ //! Uploads ray-trace data to the GPU.
+ Standard_Boolean UploadRaytraceData();
- //! Writes scene geometry to OpenCl device.
- Standard_Boolean WriteRaytraceSceneToDevice();
+ //! Resizes OpenGL frame buffers.
+ Standard_Boolean ResizeRaytraceBuffers (const Standard_Integer theSizeX,
+ const Standard_Integer theSizeY);
- //! 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);
+ //! Generates viewing rays for corners of screen quad.
+ void UpdateCamera (const NCollection_Mat4<GLdouble>& theOrientation,
+ const NCollection_Mat4<GLdouble>& theViewMapping,
+ OpenGl_Vec3 theOrigins[4],
+ OpenGl_Vec3 theDirects[4]);
- //! 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);
+ //! 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 OpenCL ray tracing.
+ //! 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;
//! 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;
-
- //! 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 texture to store rendering result.
- Handle(OpenGl_Texture) myRaytraceOutputTexture;
- //! OpenGL texture to store anti-aliasing result.
- Handle(OpenGl_Texture) myRaytraceOutputTextureAA;
-
- //! OpenCL buffer of material properties.
- cl_mem myRaytraceMaterialBuffer;
- //! OpenCL buffer of light source properties.
- cl_mem myRaytraceLightSourceBuffer;
-
- //! OpenCL buffer of vertex coords.
- cl_mem myGeometryVertexBuffer;
- //! OpenCL buffer of vertex normals.
- cl_mem myGeometryNormalBuffer;
- //! OpenCL buffer of triangle indices.
- cl_mem myGeometryTriangBuffer;
-
- //! 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;
-
- //! 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;
+ //! Actual size of traversal stack in shader program.
+ Standard_Integer myTraversalStackSize;
+
+ //! OpenGL/GLSL source of ray-tracing fragment shader.
+ ShaderSource myRaytraceShaderSource;
+ //! OpenGL/GLSL source of adaptive-AA fragment shader.
+ ShaderSource myPostFSAAShaderSource;
+
+ //! OpenGL/GLSL ray-tracing fragment shader.
+ Handle(OpenGl_ShaderObject) myRaytraceShader;
+ //! OpenGL/GLSL adaptive-AA fragment shader.
+ Handle(OpenGl_ShaderObject) myPostFSAAShader;
+
+ //! OpenGL/GLSL ray-tracing shader program.
+ Handle(OpenGl_ShaderProgram) myRaytraceProgram;
+ //! OpenGL/GLSL adaptive-AA shader program.
+ Handle(OpenGl_ShaderProgram) myPostFSAAProgram;
+
+ //! 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;
+
+ //! 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;
+
+ //! 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;
//! State of OpenGL structures reflected to ray-tracing.
std::map<const OpenGl_Structure*, Standard_Size> myStructureStates;
-#endif // HAVE_OPENCL
+ //! Cached locations of frequently used uniform variables.
+ Standard_Integer myUniformLocations[2][OpenGl_RT_NbVariables];
protected: //! @name protected fields
// Alternatively, this file may be used under the terms of Open CASCADE
// commercial license or contractual agreement.
-#ifdef HAVE_CONFIG_H
- #include <config.h>
-#endif
-
-#ifdef HAVE_OPENCL
-
-#include <OpenGl_Cl.hxx>
-
-#if defined(_WIN32)
-
- #include <windows.h>
- #include <wingdi.h>
-
- #pragma comment (lib, "DelayImp.lib")
- #pragma comment (lib, "OpenCL.lib")
-
-#elif defined(__APPLE__) && !defined(MACOSX_USE_GLX)
- #include <OpenGL/CGLCurrent.h>
-#else
- #include <GL/glx.h>
-#endif
-
-#include <OpenGl_Context.hxx>
+#include <NCollection_Mat4.hxx>
+#include <OpenGl_ArbFBO.hxx>
+#include <OpenGl_FrameBuffer.hxx>
#include <OpenGl_Texture.hxx>
+#include <OpenGl_VertexBuffer.hxx>
#include <OpenGl_View.hxx>
#include <OpenGl_Workspace.hxx>
+#include <OSD_File.hxx>
+#include <OSD_Protection.hxx>
#include <Standard_Assert.hxx>
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 <OSD_Timer.hxx>
#endif
-//! OpenCL source of ray-tracing kernels.
-extern const char THE_RAY_TRACE_OPENCL_SOURCE[];
-
// =======================================================================
// function : MatVecMult
// purpose : Multiples 4x4 matrix by 4D vector
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
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;
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);
}
// =======================================================================
{
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())
{
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;
}
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<Standard_Integer> (myRaytraceGeometry.Sources.size()) : 1;
-
- myRaytraceLightSourceBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
- aLightBufferSize * sizeof(OpenGl_RaytraceLight), NULL, &anError);
-
- if (myRaytraceGeometry.Sources.size() != 0)
+ if (myRaytraceLightSrcTexture.IsNull()) // create light source buffer
{
- const void* aDataPtr = myRaytraceGeometry.Sources.front().Packed();
-
- anError |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0,
- aLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr, 0, NULL, NULL);
- }
+ myRaytraceLightSrcTexture = new OpenGl_TextureBufferArb;
+ if (!myRaytraceLightSrcTexture->Create (myGlContext))
+ {
#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 create light source buffer" << std::endl;
#endif
-
- return Standard_True;
-}
-
-// =======================================================================
-// function : CheckOpenCL
-// purpose : Checks OpenCL dynamic library availability
-// =======================================================================
-Standard_Boolean CheckOpenCL()
-{
-#if defined ( _WIN32 )
-
- __try
- {
- cl_uint aNbPlatforms;
- clGetPlatformIDs (0, NULL, &aNbPlatforms);
+ return Standard_False;
+ }
}
- __except (EXCEPTION_EXECUTE_HANDLER)
+
+ if (myRaytraceGeometry.Sources.size() != 0)
{
- return Standard_False;
- }
+ const GLfloat* aDataPtr = myRaytraceGeometry.Sources.front().Packed();
+
+ bool aResult = myRaytraceLightSrcTexture->Init (
+ myGlContext, 4, myRaytraceGeometry.Sources.size() * 2, aDataPtr);
+ if (!aResult)
+ {
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error: Failed to upload light source buffer" << std::endl;
#endif
+ return Standard_False;
+ }
+ }
return Standard_True;
}
// =======================================================================
-// function : InitOpenCL
-// purpose : Initializes OpenCL objects
+// function : UpdateRaytraceEnvironmentMap
+// purpose : Updates environment map for ray-tracing
// =======================================================================
-Standard_Boolean OpenGl_Workspace::InitOpenCL()
+Standard_Boolean OpenGl_Workspace::UpdateRaytraceEnvironmentMap()
{
- if (myComputeInitStatus != OpenGl_CLIS_NONE)
- {
- return myComputeInitStatus == OpenGl_CLIS_INIT;
- }
-
- 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!");
+ if (myView.IsNull())
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;
- }
+ if (myViewModificationStatus == myView->ModificationState())
+ return Standard_True;
- // 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)
+ for (Standard_Integer anIdx = 0; anIdx < 2; ++anIdx)
{
- char aName[256];
- anError = clGetPlatformInfo (aPlatforms[aPlatIter], CL_PLATFORM_NAME,
- sizeof(aName), aName, NULL);
- if (anError != CL_SUCCESS)
- {
- continue;
- }
+ const Handle(OpenGl_ShaderProgram)& aProgram =
+ anIdx == 0 ? myRaytraceProgram : myPostFSAAProgram;
- if (strncmp (aName, "NVIDIA", strlen ("NVIDIA")) == 0)
+ if (!aProgram.IsNull())
{
- 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);
+ aProgram->Bind (myGlContext);
- 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);
}
- }
- }
-
- if (anError != CL_SUCCESS)
- {
- 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;
+ aProgram->SetSampler (myGlContext,
+ "uEnvironmentMapTexture", OpenGl_RT_EnvironmentMapTexture);
+ }
}
- // 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!");
+ OpenGl_ShaderProgram::Unbind (myGlContext);
- return Standard_False;
- }
+ myViewModificationStatus = myView->ModificationState();
- myComputeInitStatus = OpenGl_CLIS_INIT; // initialized in normal way
return Standard_True;
}
// =======================================================================
-// function : GetOpenClDeviceInfo
-// purpose : Returns information about device used for computations
+// function : Source
+// purpose : Returns shader source combined with prefix
// =======================================================================
-Standard_Boolean OpenGl_Workspace::GetOpenClDeviceInfo (NCollection_DataMap<TCollection_AsciiString,
- TCollection_AsciiString>& theInfo) const
+TCollection_AsciiString OpenGl_Workspace::ShaderSource::Source() const
{
- theInfo.Clear();
- if (myComputeContext == NULL)
- {
- return Standard_False;
- }
+ static const TCollection_AsciiString aVersion = "#version 140";
- 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)
+ if (myPrefix.IsEmpty())
{
- return Standard_False;
+ return aVersion + "\n" + mySource;
}
- 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;
+ return aVersion + "\n" + myPrefix + "\n" + mySource;
}
// =======================================================================
-// function : ReleaseOpenCL
-// purpose : Releases resources of OpenCL objects
+// function : Load
+// purpose : Loads shader source from specified files
// =======================================================================
-void OpenGl_Workspace::ReleaseOpenCL()
+void OpenGl_Workspace::ShaderSource::Load (
+ const TCollection_AsciiString* theFileNames, const Standard_Integer theCount)
{
- clReleaseKernel (myRaytraceRenderKernel);
- clReleaseKernel (myRaytraceSmoothKernel);
+ mySource.Clear();
- clReleaseProgram (myRaytraceProgram);
- clReleaseCommandQueue (myComputeQueue);
-
- clReleaseMemObject (myRaytraceOutputImage);
- clReleaseMemObject (myRaytraceEnvironment);
- clReleaseMemObject (myRaytraceOutputImageAA);
+ for (Standard_Integer anIndex = 0; anIndex < theCount; ++anIndex)
+ {
+ OSD_File aFile (theFileNames[anIndex]);
- clReleaseMemObject (myRaytraceMaterialBuffer);
- clReleaseMemObject (myRaytraceLightSourceBuffer);
+ Standard_ASSERT_RETURN (aFile.Exists(),
+ "Error: Failed to find shader source file", /* none */);
- clReleaseMemObject (mySceneNodeInfoBuffer);
- clReleaseMemObject (mySceneMinPointBuffer);
- clReleaseMemObject (mySceneMaxPointBuffer);
+ aFile.Open (OSD_ReadOnly, OSD_Protection());
- clReleaseMemObject (myObjectNodeInfoBuffer);
- clReleaseMemObject (myObjectMinPointBuffer);
- clReleaseMemObject (myObjectMaxPointBuffer);
+ TCollection_AsciiString aSource;
- clReleaseMemObject (myGeometryVertexBuffer);
- clReleaseMemObject (myGeometryNormalBuffer);
- clReleaseMemObject (myGeometryTriangBuffer);
+ Standard_ASSERT_RETURN (aFile.IsOpen(),
+ "Error: Failed to open shader source file", /* none */);
- clReleaseContext (myComputeContext);
+ aFile.Read (aSource, (Standard_Integer) aFile.Size());
- if (!myGlContext.IsNull())
- {
- if (!myRaytraceOutputTexture.IsNull())
- myGlContext->DelayedRelease (myRaytraceOutputTexture);
- myRaytraceOutputTexture.Nullify();
+ if (!aSource.IsEmpty())
+ {
+ mySource += TCollection_AsciiString ("\n") + aSource;
+ }
- if (!myRaytraceOutputTextureAA.IsNull())
- myGlContext->DelayedRelease (myRaytraceOutputTextureAA);
- myRaytraceOutputTextureAA.Nullify();
+ aFile.Close();
}
}
// =======================================================================
-// function : ResizeRaytraceOutputBuffer
-// purpose : Resizes OpenCL output image
+// function : LoadShader
+// purpose : Creates new shader object with specified source
// =======================================================================
-Standard_Boolean OpenGl_Workspace::ResizeRaytraceOutputBuffer (const cl_int theSizeX,
- const cl_int theSizeY)
+Handle(OpenGl_ShaderObject) OpenGl_Workspace::LoadShader (const ShaderSource& theSource, GLenum theType)
{
- if (myComputeContext == NULL)
- {
- return Standard_False;
- }
+ Handle(OpenGl_ShaderObject) aShader = new OpenGl_ShaderObject (theType);
- if (!myRaytraceOutputTexture.IsNull())
+ if (!aShader->Create (myGlContext))
{
- Standard_Boolean toResize = myRaytraceOutputTexture->SizeX() != theSizeX ||
- myRaytraceOutputTexture->SizeY() != theSizeY;
+ 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);
- if (!toResize)
- return Standard_True;
+ aShader->Release (myGlContext.operator->());
- if (!myGlContext.IsNull())
- {
- if (!myRaytraceOutputTexture.IsNull())
- myGlContext->DelayedRelease (myRaytraceOutputTexture);
- if (!myRaytraceOutputTextureAA.IsNull())
- myGlContext->DelayedRelease (myRaytraceOutputTextureAA);
- }
+ return Handle(OpenGl_ShaderObject)();
}
- myRaytraceOutputTexture = new OpenGl_Texture();
-
- myRaytraceOutputTexture->Create (myGlContext);
- myRaytraceOutputTexture->InitRectangle (myGlContext,
- theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>());
-
- myRaytraceOutputTextureAA = new OpenGl_Texture();
-
- myRaytraceOutputTextureAA->Create (myGlContext);
- myRaytraceOutputTextureAA->InitRectangle (myGlContext,
- theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>());
-
- if (myRaytraceOutputImage != NULL)
- clReleaseMemObject (myRaytraceOutputImage);
+ 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);
- if (myRaytraceOutputImageAA != NULL)
- clReleaseMemObject (myRaytraceOutputImageAA);
+ aShader->Release (myGlContext.operator->());
- cl_int anError = CL_SUCCESS;
+ return Handle(OpenGl_ShaderObject)();
+ }
- myRaytraceOutputImage = clCreateFromGLTexture2D (myComputeContext,
- CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTexture->TextureId(), &anError);
+ TCollection_AsciiString aBuildLog;
- if (anError != CL_SUCCESS)
+ 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 << "Error! Failed to create output image!" << std::endl;
+ std::cout << aBuildLog << std::endl;
#endif
- return Standard_False;
- }
- myRaytraceOutputImageAA = clCreateFromGLTexture2D (myComputeContext,
- CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTextureAA->TextureId(), &anError);
+ 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 (anError != CL_SUCCESS)
- {
#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to create anti-aliased output image!" << std::endl;
+ 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 Standard_False;
- }
- return Standard_True;
+ return aShader;
}
// =======================================================================
-// function : WriteRaytraceSceneToDevice
-// purpose : Writes scene geometry to OpenCL device
+// function : SafeFailBack
+// purpose : Performs safe exit when shaders initialization fails
// =======================================================================
-Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice()
+Standard_Boolean OpenGl_Workspace::SafeFailBack (const TCollection_ExtendedString& theMessage)
{
- if (myComputeContext == NULL)
- return Standard_False;
-
- cl_int anErrorRes = CL_SUCCESS;
+ myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+ GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, theMessage);
- if (mySceneNodeInfoBuffer != NULL)
- anErrorRes |= clReleaseMemObject (mySceneNodeInfoBuffer);
+ myComputeInitStatus = OpenGl_RT_FAIL;
- if (mySceneMinPointBuffer != NULL)
- anErrorRes |= clReleaseMemObject (mySceneMinPointBuffer);
-
- if (mySceneMaxPointBuffer != NULL)
- anErrorRes |= clReleaseMemObject (mySceneMaxPointBuffer);
-
- if (myObjectNodeInfoBuffer != NULL)
- anErrorRes |= clReleaseMemObject (myObjectNodeInfoBuffer);
+ ReleaseRaytraceResources();
+
+ return Standard_False;
+}
- if (myObjectMinPointBuffer != NULL)
- anErrorRes |= clReleaseMemObject (myObjectMinPointBuffer);
+// =======================================================================
+// function : InitRaytraceResources
+// purpose : Initializes OpenGL/GLSL shader programs
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::InitRaytraceResources()
+{
+ Standard_Boolean aToRebuildShaders = Standard_False;
- if (myObjectMaxPointBuffer != NULL)
- anErrorRes |= clReleaseMemObject (myObjectMaxPointBuffer);
+ if (myComputeInitStatus == OpenGl_RT_INIT)
+ {
+ if (!myIsRaytraceDataValid)
+ return Standard_True;
- if (myGeometryVertexBuffer != NULL)
- anErrorRes |= clReleaseMemObject (myGeometryVertexBuffer);
+ const Standard_Integer aRequiredStackSize =
+ myRaytraceGeometry.HighLevelTreeDepth() + myRaytraceGeometry.BottomLevelTreeDepth();
- if (myGeometryNormalBuffer != NULL)
- anErrorRes |= clReleaseMemObject (myGeometryNormalBuffer);
+ if (myTraversalStackSize < aRequiredStackSize)
+ {
+ myTraversalStackSize = Max (aRequiredStackSize, THE_DEFAULT_STACK_SIZE);
- if (myGeometryTriangBuffer != NULL)
- anErrorRes |= clReleaseMemObject (myGeometryTriangBuffer);
+ aToRebuildShaders = Standard_True;
+ }
+ else
+ {
+ if (aRequiredStackSize < myTraversalStackSize)
+ {
+ if (myTraversalStackSize > THE_DEFAULT_STACK_SIZE)
+ {
+ myTraversalStackSize = Max (aRequiredStackSize, THE_DEFAULT_STACK_SIZE);
- if (myRaytraceMaterialBuffer != NULL)
- anErrorRes |= clReleaseMemObject (myRaytraceMaterialBuffer);
+ aToRebuildShaders = Standard_True;
+ }
+ }
+ }
- if (anErrorRes != CL_SUCCESS)
- {
+ if (aToRebuildShaders)
+ {
#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to release OpenCL buffers" << std::endl;
+ std::cout << "Info: Rebuild shaders with stack size: " << myTraversalStackSize << std::endl;
#endif
- return Standard_False;
- }
- /////////////////////////////////////////////////////////////////////////////
- // Create material buffer
+ TCollection_AsciiString aStackSizeStr =
+ TCollection_AsciiString ("#define STACK_SIZE ") + TCollection_AsciiString (myTraversalStackSize);
- const size_t aMaterialBufferSize =
- myRaytraceGeometry.Materials.size() != 0 ? myRaytraceGeometry.Materials.size() : 1;
+ myRaytraceShaderSource.SetPrefix (aStackSizeStr);
+ myPostFSAAShaderSource.SetPrefix (aStackSizeStr);
- myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL, &anErrorRes);
+ if (!myRaytraceShader->LoadSource (myGlContext, myRaytraceShaderSource.Source())
+ || !myPostFSAAShader->LoadSource (myGlContext, myPostFSAAShaderSource.Source()))
+ {
+ return Standard_False;
+ }
- if (anErrorRes != CL_SUCCESS)
- {
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to create OpenCL material buffer" << std::endl;
-#endif
- return Standard_False;
+ if (!myRaytraceShader->Compile (myGlContext)
+ || !myPostFSAAShader->Compile (myGlContext))
+ {
+ return Standard_False;
+ }
+
+ if (!myRaytraceProgram->Link (myGlContext)
+ || !myPostFSAAProgram->Link (myGlContext))
+ {
+ return Standard_False;
+ }
+ }
}
- /////////////////////////////////////////////////////////////////////////////
- // Create BVHs buffers
+ if (myComputeInitStatus == OpenGl_RT_NONE)
+ {
+ if (!myGlContext->IsGlGreaterEqual (3, 1))
+ {
+ const TCollection_ExtendedString aMessage = "Ray-tracing requires OpenGL 3.1 and higher";
- cl_int anErrorTmp = CL_SUCCESS;
+ myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
+ GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
- const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = myRaytraceGeometry.BVH();
+ return Standard_False;
+ }
- const size_t aSceneMinPointBufferSize =
- aBVH->MinPointBuffer().size() != 0 ? aBVH->MinPointBuffer().size() : 1;
+ TCollection_AsciiString aFolder = Graphic3d_ShaderProgram::ShadersFolder();
- mySceneMinPointBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aSceneMinPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp);
- anErrorRes |= anErrorTmp;
+ 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;
+ }
- const size_t aSceneMaxPointBufferSize =
- aBVH->MaxPointBuffer().size() != 0 ? aBVH->MaxPointBuffer().size() : 1;
+ if (myIsRaytraceDataValid)
+ {
+ myTraversalStackSize = Max (THE_DEFAULT_STACK_SIZE,
+ myRaytraceGeometry.HighLevelTreeDepth() + myRaytraceGeometry.BottomLevelTreeDepth());
+ }
- mySceneMaxPointBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aSceneMaxPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp);
- anErrorRes |= anErrorTmp;
+ {
+ Handle(OpenGl_ShaderObject) aBasicVertShader = LoadShader (
+ ShaderSource (aFolder + "/RaytraceBase.vs"), GL_VERTEX_SHADER);
- const size_t aSceneNodeInfoBufferSize =
- aBVH->NodeInfoBuffer().size() != 0 ? aBVH->NodeInfoBuffer().size() : 1;
+ if (aBasicVertShader.IsNull())
+ {
+ return SafeFailBack ("Failed to set vertex shader source");
+ }
- mySceneNodeInfoBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aSceneNodeInfoBufferSize * sizeof(cl_int4), NULL, &anErrorTmp);
- anErrorRes |= anErrorTmp;
+ TCollection_AsciiString aFiles[] = { aFolder + "/RaytraceBase.fs", aFolder + "/RaytraceRender.fs" };
- if (anErrorRes != CL_SUCCESS)
- {
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to create OpenCL buffers for high-level scene BVH" << std::endl;
-#endif
- return Standard_False;
- }
+ myRaytraceShaderSource.Load (aFiles, 2);
- Standard_Integer aTotalVerticesNb = 0;
- Standard_Integer aTotalElementsNb = 0;
- Standard_Integer aTotalBVHNodesNb = 0;
+ TCollection_AsciiString aStackSizeStr =
+ TCollection_AsciiString ("#define STACK_SIZE ") + TCollection_AsciiString (myTraversalStackSize);
- for (Standard_Integer anElemIndex = 0; anElemIndex < myRaytraceGeometry.Size(); ++anElemIndex)
- {
- OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
- myRaytraceGeometry.Objects().ChangeValue (anElemIndex).operator->());
+ myRaytraceShaderSource.SetPrefix (aStackSizeStr);
- Standard_ASSERT_RETURN (aTriangleSet != NULL,
- "Error! Failed to get triangulation of OpenGL element", Standard_False);
+ myRaytraceShader = LoadShader (myRaytraceShaderSource, GL_FRAGMENT_SHADER);
- aTotalVerticesNb += (int)aTriangleSet->Vertices.size();
- aTotalElementsNb += (int)aTriangleSet->Elements.size();
+ if (myRaytraceShader.IsNull())
+ {
+ aBasicVertShader->Release (myGlContext.operator->());
- Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(),
- "Error! Failed to get bottom-level BVH of OpenGL element", Standard_False);
+ return SafeFailBack ("Failed to set ray-trace fragment shader source");
+ }
- aTotalBVHNodesNb += (int)aTriangleSet->BVH()->NodeInfoBuffer().size();
- }
+ myRaytraceProgram = new OpenGl_ShaderProgram;
- aTotalBVHNodesNb = aTotalBVHNodesNb > 0 ? aTotalBVHNodesNb : 1;
+ if (!myRaytraceProgram->Create (myGlContext))
+ {
+ aBasicVertShader->Release (myGlContext.operator->());
- myObjectNodeInfoBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_int4), NULL, &anErrorTmp);
- anErrorRes |= anErrorTmp;
+ return SafeFailBack ("Failed to create ray-trace shader program");
+ }
- myObjectMinPointBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp);
- anErrorRes |= anErrorTmp;
+ if (!myRaytraceProgram->AttachShader (myGlContext, aBasicVertShader)
+ || !myRaytraceProgram->AttachShader (myGlContext, myRaytraceShader))
+ {
+ aBasicVertShader->Release (myGlContext.operator->());
- myObjectMaxPointBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp);
- anErrorRes |= anErrorTmp;
+ return SafeFailBack ("Failed to attach ray-trace shader objects");
+ }
- if (anErrorRes != CL_SUCCESS)
- {
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to create OpenCL buffers for bottom-level scene BVHs" << std::endl;
-#endif
- return Standard_False;
- }
+ if (!myRaytraceProgram->Link (myGlContext))
+ {
+ TCollection_AsciiString aLinkLog;
- /////////////////////////////////////////////////////////////////////////////
- // Create geometry buffers
+ if (myRaytraceProgram->FetchInfoLog (myGlContext, aLinkLog))
+ {
+ #ifdef RAY_TRACE_PRINT_INFO
+ std::cout << aLinkLog << std::endl;
+ #endif
+ }
- aTotalVerticesNb = aTotalVerticesNb > 0 ? aTotalVerticesNb : 1;
+ return SafeFailBack ("Failed to link ray-trace shader program");
+ }
+ }
- myGeometryVertexBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp);
- anErrorRes |= anErrorTmp;
+ {
+ Handle(OpenGl_ShaderObject) aBasicVertShader = LoadShader (
+ ShaderSource (aFolder + "/RaytraceBase.vs"), GL_VERTEX_SHADER);
- myGeometryNormalBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp);
- anErrorRes |= anErrorTmp;
+ if (aBasicVertShader.IsNull())
+ {
+ return SafeFailBack ("Failed to set vertex shader source");
+ }
- aTotalElementsNb = aTotalElementsNb > 0 ? aTotalElementsNb : 1;
+ TCollection_AsciiString aFiles[] = { aFolder + "/RaytraceBase.fs", aFolder + "/RaytraceSmooth.fs" };
- myGeometryTriangBuffer = clCreateBuffer (myComputeContext,
- CL_MEM_READ_ONLY, aTotalElementsNb * sizeof(cl_int4), NULL, &anErrorTmp);
- anErrorRes |= anErrorTmp;
+ myPostFSAAShaderSource.Load (aFiles, 2);
- if (anErrorRes != CL_SUCCESS)
- {
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to create OpenCL geometry buffers" << std::endl;
-#endif
- return Standard_False;
- }
+ TCollection_AsciiString aStackSizeStr =
+ TCollection_AsciiString ("#define STACK_SIZE ") + TCollection_AsciiString (myTraversalStackSize);
- /////////////////////////////////////////////////////////////////////////////
- // Write BVH and geometry buffers
+ myPostFSAAShaderSource.SetPrefix (aStackSizeStr);
+
+ myPostFSAAShader = LoadShader (myPostFSAAShaderSource, GL_FRAGMENT_SHADER);
- if (aBVH->NodeInfoBuffer().size() != 0)
- {
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneNodeInfoBuffer, CL_FALSE, 0,
- aSceneNodeInfoBufferSize * sizeof(cl_int4), &aBVH->NodeInfoBuffer().front(), 0, NULL, NULL);
+ if (myPostFSAAShader.IsNull())
+ {
+ aBasicVertShader->Release (myGlContext.operator->());
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMinPointBuffer, CL_FALSE, 0,
- aSceneMinPointBufferSize * sizeof(cl_float4), &aBVH->MinPointBuffer().front(), 0, NULL, NULL);
+ return SafeFailBack ("Failed to set FSAA fragment shader source");
+ }
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMaxPointBuffer, CL_FALSE, 0,
- aSceneMaxPointBufferSize * sizeof(cl_float4), &aBVH->MaxPointBuffer().front(), 0, NULL, NULL);
+ myPostFSAAProgram = new OpenGl_ShaderProgram;
- anErrorRes |= clFinish (myComputeQueue);
+ if (!myPostFSAAProgram->Create (myGlContext))
+ {
+ aBasicVertShader->Release (myGlContext.operator->());
- if (anErrorRes != CL_SUCCESS)
- {
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to write OpenCL buffers for high-level scene BVH" << std::endl;
-#endif
- return Standard_False;
+ 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");
+ }
}
+ }
- for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx)
+ if (myComputeInitStatus == OpenGl_RT_NONE || aToRebuildShaders)
+ {
+ for (Standard_Integer anIndex = 0; anIndex < 2; ++anIndex)
{
- if (!aBVH->IsOuter (aNodeIdx))
- continue;
-
- OpenGl_TriangleSet* aTriangleSet = myRaytraceGeometry.TriangleSet (aNodeIdx);
+ 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);
+ }
- Standard_ASSERT_RETURN (aTriangleSet != NULL,
- "Error! Failed to get triangulation of OpenGL element", Standard_False);
+ 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");
+ }
- const size_t aBVHBuffserSize =
- aTriangleSet->BVH()->NodeInfoBuffer().size() != 0 ? aTriangleSet->BVH()->NodeInfoBuffer().size() : 1;
+ OpenGl_ShaderProgram::Unbind (myGlContext);
+ }
- const Standard_Integer aBVHOffset = myRaytraceGeometry.AccelerationOffset (aNodeIdx);
+ if (myComputeInitStatus != OpenGl_RT_NONE)
+ {
+ return myComputeInitStatus == OpenGl_RT_INIT;
+ }
- Standard_ASSERT_RETURN (aBVHOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
- "Error! Failed to get offset for bottom-level BVH", Standard_False);
+ if (myRaytraceFBO1.IsNull())
+ {
+ myRaytraceFBO1 = new OpenGl_FrameBuffer;
+ }
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectNodeInfoBuffer, CL_FALSE, aBVHOffset * sizeof(cl_int4),
- aBVHBuffserSize * sizeof(cl_int4), &aTriangleSet->BVH()->NodeInfoBuffer().front(), 0, NULL, NULL);
+ if (myRaytraceFBO2.IsNull())
+ {
+ myRaytraceFBO2 = new OpenGl_FrameBuffer;
+ }
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMinPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4),
- aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MinPointBuffer().front(), 0, NULL, NULL);
+ 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 };
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMaxPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4),
- aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MaxPointBuffer().front(), 0, NULL, NULL);
+ myRaytraceScreenQuad.Init (myGlContext, 3, 6, aVertices);
- anErrorRes |= clFinish (myComputeQueue);
+ myComputeInitStatus = OpenGl_RT_INIT; // initialized in normal way
+
+ return Standard_True;
+}
- 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;
- }
+// =======================================================================
+// function : NullifyResource
+// purpose :
+// =======================================================================
+inline void NullifyResource (const Handle(OpenGl_Context)& theContext,
+ Handle(OpenGl_Resource)& theResource)
+{
+ if (!theResource.IsNull())
+ {
+ theResource->Release (theContext.operator->());
+ theResource.Nullify();
+ }
+}
- const Standard_Integer aVerticesOffset = myRaytraceGeometry.VerticesOffset (aNodeIdx);
+// =======================================================================
+// function : ReleaseRaytraceResources
+// purpose : Releases OpenGL/GLSL shader programs
+// =======================================================================
+void OpenGl_Workspace::ReleaseRaytraceResources()
+{
+ NullifyResource (myGlContext, myRaytraceFBO1);
+ NullifyResource (myGlContext, myRaytraceFBO2);
- Standard_ASSERT_RETURN (aVerticesOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
- "Error! Failed to get offset for triangulation vertices of OpenGL element", Standard_False);
+ NullifyResource (myGlContext, myRaytraceShader);
+ NullifyResource (myGlContext, myPostFSAAShader);
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryVertexBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4),
- aTriangleSet->Vertices.size() * sizeof(cl_float4), &aTriangleSet->Vertices.front(), 0, NULL, NULL);
+ NullifyResource (myGlContext, myRaytraceProgram);
+ NullifyResource (myGlContext, myPostFSAAProgram);
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryNormalBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4),
- aTriangleSet->Normals.size() * sizeof(cl_float4), &aTriangleSet->Normals.front(), 0, NULL, NULL);
+ NullifyResource (myGlContext, mySceneNodeInfoTexture);
+ NullifyResource (myGlContext, mySceneMinPointTexture);
+ NullifyResource (myGlContext, mySceneMaxPointTexture);
- const Standard_Integer anElementsOffset = myRaytraceGeometry.ElementsOffset (aNodeIdx);
+ NullifyResource (myGlContext, myObjectNodeInfoTexture);
+ NullifyResource (myGlContext, myObjectMinPointTexture);
+ NullifyResource (myGlContext, myObjectMaxPointTexture);
- Standard_ASSERT_RETURN (anElementsOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
- "Error! Failed to get offset for triangulation elements of OpenGL element", Standard_False);
+ NullifyResource (myGlContext, myGeometryVertexTexture);
+ NullifyResource (myGlContext, myGeometryNormalTexture);
+ NullifyResource (myGlContext, myGeometryTriangTexture);
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryTriangBuffer, CL_FALSE, anElementsOffset * sizeof(cl_int4),
- aTriangleSet->Elements.size() * sizeof(cl_int4), &aTriangleSet->Elements.front(), 0, NULL, NULL);
+ NullifyResource (myGlContext, myRaytraceLightSrcTexture);
+ NullifyResource (myGlContext, myRaytraceMaterialTexture);
- anErrorRes |= clFinish (myComputeQueue);
+ if (myRaytraceScreenQuad.IsValid())
+ myRaytraceScreenQuad.Release (myGlContext.operator->());
+}
- 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 write OpenCL triangulation buffers for OpenGL element" << std::endl;
+ std::cout << "Error: OpenGL version is less than 3.1" << std::endl;
#endif
- return Standard_False;
- }
- }
+ return Standard_False;
}
/////////////////////////////////////////////////////////////////////////////
- // Write material buffer
+ // Create OpenGL texture buffers
- if (myRaytraceGeometry.Materials.size() != 0)
+ if (mySceneNodeInfoTexture.IsNull()) // create hight-level BVH buffers
{
- const void* aDataPtr = myRaytraceGeometry.Materials.front().Packed();
-
- anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceMaterialBuffer,
- CL_FALSE, 0, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr, 0, NULL, NULL);
+ mySceneNodeInfoTexture = new OpenGl_TextureBufferArb;
+ mySceneMinPointTexture = new OpenGl_TextureBufferArb;
+ mySceneMaxPointTexture = new OpenGl_TextureBufferArb;
- if (anErrorRes != CL_SUCCESS)
+ if (!mySceneNodeInfoTexture->Create (myGlContext)
+ || !mySceneMinPointTexture->Create (myGlContext)
+ || !mySceneMaxPointTexture->Create (myGlContext))
{
- #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 create buffers for high-level scene BVH" << std::endl;
+#endif
return Standard_False;
}
}
- anErrorRes |= clFinish (myComputeQueue);
-
- if (anErrorRes == CL_SUCCESS)
+ if (myObjectNodeInfoTexture.IsNull()) // create bottom-level BVH buffers
{
- myIsRaytraceDataValid = myRaytraceGeometry.Objects().Size() != 0;
- }
-#ifdef RAY_TRACE_PRINT_INFO
- else
- {
- std::cout << "Error! Failed to set scene data buffers" << std::endl;
- }
-#endif
+ 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;
+ }
+ }
- Standard_ShortReal aMemUsed = 0.f;
-
- for (Standard_Integer anElemIdx = 0; anElemIdx < myRaytraceGeometry.Size(); ++anElemIdx)
+ if (myGeometryVertexTexture.IsNull()) // create geometry buffers
{
- OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
- myRaytraceGeometry.Objects().ChangeValue (anElemIdx).operator->());
+ myGeometryVertexTexture = new OpenGl_TextureBufferArb;
+ myGeometryNormalTexture = new OpenGl_TextureBufferArb;
+ myGeometryTriangTexture = new OpenGl_TextureBufferArb;
- aMemUsed += static_cast<Standard_ShortReal> (
- aTriangleSet->Vertices.size() * sizeof (BVH_Vec4f));
- aMemUsed += static_cast<Standard_ShortReal> (
- aTriangleSet->Normals.size() * sizeof (BVH_Vec4f));
- aMemUsed += static_cast<Standard_ShortReal> (
- aTriangleSet->Elements.size() * sizeof (BVH_Vec4i));
-
- aMemUsed += static_cast<Standard_ShortReal> (
- aTriangleSet->BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
- aMemUsed += static_cast<Standard_ShortReal> (
- aTriangleSet->BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
- aMemUsed += static_cast<Standard_ShortReal> (
- aTriangleSet->BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
+ 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;
+ }
}
- aMemUsed += static_cast<Standard_ShortReal> (
- myRaytraceGeometry.BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
- aMemUsed += static_cast<Standard_ShortReal> (
- myRaytraceGeometry.BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
- aMemUsed += static_cast<Standard_ShortReal> (
- myRaytraceGeometry.BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
-
- std::cout << "GPU Memory Used (MB): ~" << aMemUsed / 1048576 << std::endl;
+ 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;
+ }
+ }
- return (CL_SUCCESS == anErrorRes);
-}
+ /////////////////////////////////////////////////////////////////////////////
+ // Write OpenGL texture buffers
-// Use it to estimate the optimal size of OpenCL work group
-// #define OPENCL_GROUP_SIZE_TEST
+ const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = myRaytraceGeometry.BVH();
-// =======================================================================
-// function : RunRaytraceOpenCLKernels
-// purpose : Runs OpenCL ray-tracing kernels
-// =======================================================================
-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)
-{
- if (myRaytraceRenderKernel == NULL || myComputeQueue == NULL)
- return Standard_False;
+ bool aResult = true;
- ////////////////////////////////////////////////////////////////////////
- // 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)
- {
- const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of ray-tracing kernel!";
+ if (!aBVH->NodeInfoBuffer().empty())
+ {
+ aResult &= mySceneNodeInfoTexture->Init (myGlContext, 4,
+ aBVH->NodeInfoBuffer().size(), reinterpret_cast<const GLuint*> (&aBVH->NodeInfoBuffer().front()));
- myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
- GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+ aResult &= mySceneMinPointTexture->Init (myGlContext, 4,
+ aBVH->MinPointBuffer().size(), reinterpret_cast<const GLfloat*> (&aBVH->MinPointBuffer().front()));
+
+ aResult &= mySceneMaxPointTexture->Init (myGlContext, 4,
+ aBVH->MaxPointBuffer().size(), reinterpret_cast<const GLfloat*> (&aBVH->MaxPointBuffer().front()));
+ }
+ if (!aResult)
+ {
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error: Failed to upload buffers for high-level scene BVH" << std::endl;
+#endif
return Standard_False;
}
- // 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!";
+ Standard_Size aTotalVerticesNb = 0;
+ Standard_Size aTotalElementsNb = 0;
+ Standard_Size aTotalBVHNodesNb = 0;
- myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
- GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+ for (Standard_Integer anElemIndex = 0; anElemIndex < myRaytraceGeometry.Size(); ++anElemIndex)
+ {
+ OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
+ myRaytraceGeometry.Objects().ChangeValue (anElemIndex).operator->());
- return Standard_False;
- }
- }
+ Standard_ASSERT_RETURN (aTriangleSet != NULL,
+ "Error: Failed to get triangulation of OpenGL element", Standard_False);
+
+ aTotalVerticesNb += aTriangleSet->Vertices.size();
+ aTotalElementsNb += aTriangleSet->Elements.size();
- ////////////////////////////////////////////////////////////////////////
- // Set work size
+ Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(),
+ "Error: Failed to get bottom-level BVH of OpenGL element", Standard_False);
- size_t aLocWorkSize[] = { myIsAmdComputePlatform ? 2 : 4, 32 };
+ aTotalBVHNodesNb += aTriangleSet->BVH()->NodeInfoBuffer().size();
+ }
-#ifdef OPENCL_GROUP_SIZE_TEST
- for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1)
- for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1)
+ if (aTotalBVHNodesNb != 0)
{
- aLocWorkSize[0] = aLocX;
- aLocWorkSize[1] = aLocY;
-#endif
+ aResult &= myObjectNodeInfoTexture->Init (
+ myGlContext, 4, aTotalBVHNodesNb, static_cast<const GLuint*> (NULL));
- size_t aWorkSizeX = theSizeX;
- if (aWorkSizeX % aLocWorkSize[0] != 0)
- aWorkSizeX += aLocWorkSize[0] - aWorkSizeX % aLocWorkSize[0];
+ aResult &= myObjectMinPointTexture->Init (
+ myGlContext, 4, aTotalBVHNodesNb, static_cast<const GLfloat*> (NULL));
- size_t aWokrSizeY = theSizeY;
- if (aWokrSizeY % aLocWorkSize[1] != 0 )
- aWokrSizeY += aLocWorkSize[1] - aWokrSizeY % aLocWorkSize[1];
+ aResult &= myObjectMaxPointTexture->Init (
+ myGlContext, 4, aTotalBVHNodesNb, static_cast<const GLfloat*> (NULL));
+ }
- size_t aTotWorkSize[] = { aWorkSizeX, aWokrSizeY };
+ if (!aResult)
+ {
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error: Failed to upload buffers for bottom-level scene BVH" << std::endl;
+#endif
+ return Standard_False;
+ }
- cl_event anEvent = NULL, anEventSmooth = NULL;
+ if (aTotalElementsNb != 0)
+ {
+ aResult &= myGeometryTriangTexture->Init (
+ myGlContext, 4, aTotalElementsNb, static_cast<const GLuint*> (NULL));
+ }
- anError = clEnqueueNDRangeKernel (myComputeQueue,
- myRaytraceRenderKernel, 2, NULL, aTotWorkSize, aLocWorkSize, 0, NULL, &anEvent);
+ if (aTotalVerticesNb != 0)
+ {
+ aResult &= myGeometryVertexTexture->Init (
+ myGlContext, 4, aTotalVerticesNb, static_cast<const GLfloat*> (NULL));
- if (anError != CL_SUCCESS)
- {
- const TCollection_ExtendedString aMessage = "Error! Failed to execute the ray-tracing kernel!";
+ aResult &= myGeometryNormalTexture->Init (
+ myGlContext, 4, aTotalVerticesNb, static_cast<const GLfloat*> (NULL));
+ }
- myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
- GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+ if (!aResult)
+ {
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error: Failed to upload buffers for scene geometry" << std::endl;
+#endif
+ return Standard_False;
+ }
- return Standard_False;
- }
+ for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx)
+ {
+ if (!aBVH->IsOuter (aNodeIdx))
+ continue;
- clWaitForEvents (1, &anEvent);
+ OpenGl_TriangleSet* aTriangleSet = myRaytraceGeometry.TriangleSet (aNodeIdx);
- if (theCView.IsAntialiasingEnabled)
- {
- size_t aLocWorkSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4,
- myIsAmdComputePlatform ? 8 : 32 };
+ Standard_ASSERT_RETURN (aTriangleSet != NULL,
+ "Error: Failed to get triangulation of OpenGL element", Standard_False);
-#ifdef OPENCL_GROUP_SIZE_TEST
- aLocWorkSizeSmooth[0] = aLocX;
- aLocWorkSizeSmooth[1] = aLocY;
-#endif
+ const Standard_Integer aBVHOffset = myRaytraceGeometry.AccelerationOffset (aNodeIdx);
- aWorkSizeX = theSizeX;
- if (aWorkSizeX % aLocWorkSizeSmooth[0] != 0)
- aWorkSizeX += aLocWorkSizeSmooth[0] - aWorkSizeX % aLocWorkSizeSmooth[0];
+ Standard_ASSERT_RETURN (aBVHOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
+ "Error: Failed to get offset for bottom-level BVH", Standard_False);
- size_t aWokrSizeY = theSizeY;
- if (aWokrSizeY % aLocWorkSizeSmooth[1] != 0 )
- aWokrSizeY += aLocWorkSizeSmooth[1] - aWokrSizeY % aLocWorkSizeSmooth[1];
+ const size_t aBVHBuffserSize = aTriangleSet->BVH()->NodeInfoBuffer().size();
- size_t aTotWorkSizeSmooth [] = { aWorkSizeX, aWokrSizeY };
+ if (aBVHBuffserSize != 0)
+ {
+ aResult &= myObjectNodeInfoTexture->SubData (myGlContext, aBVHOffset,
+ aBVHBuffserSize, reinterpret_cast<const GLuint*> (&aTriangleSet->BVH()->NodeInfoBuffer().front()));
- anError = clEnqueueNDRangeKernel (myComputeQueue, myRaytraceSmoothKernel,
- 2, NULL, aTotWorkSizeSmooth, aLocWorkSizeSmooth, 0, NULL, &anEventSmooth);
+ aResult &= myObjectMinPointTexture->SubData (myGlContext, aBVHOffset,
+ aBVHBuffserSize, reinterpret_cast<const GLfloat*> (&aTriangleSet->BVH()->MinPointBuffer().front()));
- clWaitForEvents (1, &anEventSmooth);
+ aResult &= myObjectMaxPointTexture->SubData (myGlContext, aBVHOffset,
+ aBVHBuffserSize, reinterpret_cast<const GLfloat*> (&aTriangleSet->BVH()->MaxPointBuffer().front()));
- if (anError != CL_SUCCESS)
+ if (!aResult)
{
- 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);
-
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error: Failed to upload buffers for bottom-level scene BVHs" << std::endl;
+#endif
return Standard_False;
}
}
-#if defined (RAY_TRACE_PRINT_INFO) || defined (OPENCL_GROUP_SIZE_TEST)
+ const Standard_Integer aVerticesOffset = myRaytraceGeometry.VerticesOffset (aNodeIdx);
- static cl_ulong ttt1 = 10000000000;
- static cl_ulong ttt2 = 10000000000;
+ Standard_ASSERT_RETURN (aVerticesOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
+ "Error: Failed to get offset for triangulation vertices of OpenGL element", Standard_False);
- cl_ulong aBegTime = 0;
- cl_ulong aEndTime = 0;
+ if (!aTriangleSet->Vertices.empty())
+ {
+ aResult &= myGeometryNormalTexture->SubData (myGlContext, aVerticesOffset,
+ aTriangleSet->Normals.size(), reinterpret_cast<const GLfloat*> (&aTriangleSet->Normals.front()));
- clGetEventProfilingInfo (anEvent,
- CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL);
- clGetEventProfilingInfo (anEvent,
- CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL);
+ aResult &= myGeometryVertexTexture->SubData (myGlContext, aVerticesOffset,
+ aTriangleSet->Vertices.size(), reinterpret_cast<const GLfloat*> (&aTriangleSet->Vertices.front()));
+ }
- ttt1 = aEndTime - aBegTime < ttt1 ? aEndTime - aBegTime : ttt1;
+ const Standard_Integer anElementsOffset = myRaytraceGeometry.ElementsOffset (aNodeIdx);
- std::cout << "\tRender time (ms): " << ttt1 / 1e6f << std::endl;
+ Standard_ASSERT_RETURN (anElementsOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
+ "Error: Failed to get offset for triangulation elements of OpenGL element", Standard_False);
- if (theCView.IsAntialiasingEnabled)
+ if (!aTriangleSet->Elements.empty())
{
- 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;
+ aResult &= myGeometryTriangTexture->SubData (myGlContext, anElementsOffset,
+ aTriangleSet->Elements.size(), reinterpret_cast<const GLuint*> (&aTriangleSet->Elements.front()));
}
+ if (!aResult)
+ {
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error: Failed to upload triangulation buffers for OpenGL element" << std::endl;
#endif
-
- if (anEvent != NULL)
- clReleaseEvent (anEvent);
-
- if (anEventSmooth != NULL)
- clReleaseEvent (anEventSmooth);
-
-#ifdef OPENCL_GROUP_SIZE_TEST
+ return Standard_False;
+ }
}
-#endif
-
- return Standard_True;
-}
-
-// =======================================================================
-// function : ComputeInverseMatrix
-// purpose : Computes inversion of 4x4 floating-point matrix
-// =======================================================================
-template <typename T>
-void ComputeInverseMatrix (const T m[16], T inv[16])
-{
- 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]);
- 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]);
-
- 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]);
+ if (myRaytraceGeometry.Materials.size() != 0)
+ {
+ const GLfloat* aDataPtr = myRaytraceGeometry.Materials.front().Packed();
- 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]);
+ aResult &= myRaytraceMaterialTexture->Init (
+ myGlContext, 4, myRaytraceGeometry.Materials.size() * 7, aDataPtr);
- 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]);
+ if (!aResult)
+ {
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error: Failed to upload material buffer" << std::endl;
+#endif
+ return Standard_False;
+ }
+ }
- 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]);
+ myIsRaytraceDataValid = myRaytraceGeometry.Objects().Size() != 0;
- 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]);
+#ifdef RAY_TRACE_PRINT_INFO
- 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]);
+ Standard_ShortReal aMemUsed = 0.f;
- 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]);
+ for (Standard_Integer anElemIdx = 0; anElemIdx < myRaytraceGeometry.Size(); ++anElemIdx)
+ {
+ OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
+ myRaytraceGeometry.Objects().ChangeValue (anElemIdx).operator->());
- 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]);
+ aMemUsed += static_cast<Standard_ShortReal> (
+ aTriangleSet->Vertices.size() * sizeof (BVH_Vec4f));
+ aMemUsed += static_cast<Standard_ShortReal> (
+ aTriangleSet->Normals.size() * sizeof (BVH_Vec4f));
+ aMemUsed += static_cast<Standard_ShortReal> (
+ aTriangleSet->Elements.size() * sizeof (BVH_Vec4i));
- 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]);
+ aMemUsed += static_cast<Standard_ShortReal> (
+ aTriangleSet->BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
+ aMemUsed += static_cast<Standard_ShortReal> (
+ aTriangleSet->BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
+ aMemUsed += static_cast<Standard_ShortReal> (
+ aTriangleSet->BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
+ }
- 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]);
+ aMemUsed += static_cast<Standard_ShortReal> (
+ myRaytraceGeometry.BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i));
+ aMemUsed += static_cast<Standard_ShortReal> (
+ myRaytraceGeometry.BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f));
+ aMemUsed += static_cast<Standard_ShortReal> (
+ myRaytraceGeometry.BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f));
- 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]);
+ std::cout << "GPU Memory Used (MB): ~" << aMemUsed / 1048576 << std::endl;
- T det = m[0] * inv[ 0] +
- m[1] * inv[ 4] +
- m[2] * inv[ 8] +
- m[3] * inv[12];
+#endif
- if (det == T (0.0)) return;
+ return aResult;
+}
- det = T (1.0) / det;
+// =======================================================================
+// function : ResizeRaytraceBuffers
+// purpose : Resizes OpenGL frame buffers
+// =======================================================================
+Standard_Boolean OpenGl_Workspace::ResizeRaytraceBuffers (const Standard_Integer theSizeX,
+ const Standard_Integer theSizeY)
+{
+ if (myRaytraceFBO1->GetVPSizeX() != theSizeX
+ || myRaytraceFBO1->GetVPSizeY() != theSizeY)
+ {
+ myRaytraceFBO1->Init (myGlContext, theSizeX, theSizeY);
+ myRaytraceFBO2->Init (myGlContext, theSizeX, theSizeY);
+ }
- for (Standard_Integer i = 0; i < 16; ++i)
- inv[i] *= det;
+ return Standard_True;
}
// =======================================================================
-// function : GenerateCornerRays
-// purpose : Generates primary rays for corners of screen quad
+// function : UpdateCamera
+// purpose : Generates viewing rays for corners of screen quad
// =======================================================================
-void GenerateCornerRays (const GLdouble theInvModelProj[16],
- cl_float theOrigins[16],
- cl_float theDirects[16])
+void OpenGl_Workspace::UpdateCamera (const NCollection_Mat4<GLdouble>& theOrientation,
+ const NCollection_Mat4<GLdouble>& theViewMapping,
+ OpenGl_Vec3 theOrigins[4],
+ OpenGl_Vec3 theDirects[4])
{
+ NCollection_Mat4<GLdouble> aInvModelProj;
+
+ // compute invserse model-view-projection matrix
+ (theViewMapping * theOrientation).Inverted (aInvModelProj);
+
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<GLfloat> (aOrigin.x());
- theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.y());
- theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.z());
- theOrigins [aOriginIndex++] = 1.f;
+ theOrigins[aOriginIndex++] = OpenGl_Vec3 (static_cast<GLfloat> (aOrigin.x()),
+ static_cast<GLfloat> (aOrigin.y()),
+ static_cast<GLfloat> (aOrigin.z()));
+
+ theDirects[aDirectIndex++] = OpenGl_Vec3 (static_cast<GLfloat> (aDirect.x() * aInvLen),
+ static_cast<GLfloat> (aDirect.y() * aInvLen),
+ static_cast<GLfloat> (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<Standard_Integer> (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);
- theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.x() * aInvLen);
- theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.y() * aInvLen);
- theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.z() * aInvLen);
- theDirects [aDirectIndex++] = 0.f;
+ 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 (!UpdateRaytraceEnvironmentMap())
+ if (!ResizeRaytraceBuffers (theSizeX, theSizeY))
return Standard_False;
- if (!UpdateRaytraceGeometry (Standard_True))
+ if (!UpdateRaytraceEnvironmentMap())
return Standard_False;
// Get model-view and projection matrices
myView->GetMatrices (theOrientation, theViewMapping);
- GLdouble aOrientationMatrix[16];
- GLdouble aViewMappingMatrix[16];
- GLdouble aOrientationInvers[16];
+ NCollection_Mat4<GLdouble> aOrientationMatrix;
+ NCollection_Mat4<GLdouble> 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<GLdouble> 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);
-
- glLoadMatrixd (aViewMappingMatrix);
- glMultMatrixd (aOrientationMatrix);
-
- GLdouble aModelProject[16];
- GLdouble aInvModelProj[16];
-
- glGetDoublev (GL_MODELVIEW_MATRIX, aModelProject);
-
- ComputeInverseMatrix (aModelProject, aInvModelProj);
+ OpenGl_Vec3 aOrigins[4];
+ OpenGl_Vec3 aDirects[4];
- 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 |
glClear (GL_COLOR_BUFFER_BIT);
- Handle(OpenGl_Workspace) aWorkspace (this);
- myView->DrawBackground (aWorkspace);
-
- // Draw dummy quad to show result image
- glEnable (GL_COLOR_MATERIAL);
- glEnable (GL_BLEND);
+ if (theFrameBuffer != NULL)
+ theFrameBuffer->BindBuffer (myGlContext);
- 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
myBackBufferRestored = Standard_False;
}
else
+ {
glFlush();
+ }
return Standard_True;
}
-
-#endif
--- /dev/null
+//! 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);
+}
--- /dev/null
+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
--- /dev/null
+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
--- /dev/null
+//! 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
CSF_objc
CSF_Appkit
CSF_IOKit
-CSF_OPENCL
CSF_FREETYPE
CSF_GL2PS
CSF_user32
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<Graphic3d_CView*> (ViewerTest::CurrentView()->View()->CView());
- NCollection_DataMap<TCollection_AsciiString, TCollection_AsciiString> 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<TCollection_AsciiString, TCollection_AsciiString>::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
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);
# activate ray-tracing
vraytrace 1
-vclinfo
set aModeNum 0
for { set aAAMode 0 } { $aAAMode <= 1 } { incr aAAMode } {
# turn on ray tracing
vraytrace 1
-vclinfo
vdump $::imagedir/${::casename}_rt1.png
vclear
}
vraytrace 1
-vclinfo
set aColorL [vreadpixel 150 250 rgb name]
set aColorR [vreadpixel 250 250 rgb name]
#if { "$aColorL" != "GREEN3" || "$aColorR" != "GREEN4" } {
projects / occt.git / commitdiff