Updated test case v3d/raytrace/bug24130. The new version checks correctness of shadows.
return;
}
- BVH_Box<T, N> aBox;
- for (Standard_Integer anIndex = 0; anIndex < BVH_ObjectSet<T, N>::myObjects.Size(); ++anIndex)
- {
- aBox.Combine (BVH_ObjectSet<T, N>::myObjects.Value (anIndex)->Box());
- }
+ myBuilder->Build (this, myBVH.operator->(), Box());
- myBuilder->Build (this, myBVH.operator->(), aBox);
myIsDirty = Standard_False;
}
Handle_OpenGl_ShaderProgram.hxx
Handle_OpenGl_ShaderManager.hxx
OpenGl_Cl.hxx
-OpenGl_AABB.hxx
-OpenGl_AABB.cxx
OpenGl_SceneGeometry.hxx
OpenGl_SceneGeometry.cxx
-OpenGl_RaytraceTypes.hxx
OpenGl_RaytraceSource.cxx
OpenGl_Workspace_Raytrace.cxx
OpenGl_Flipper.hxx
+++ /dev/null
-// Created on: 2013-08-27
-// Created by: Denis BOGOLEPOV
-// Copyright (c) 2013 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 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
-
-#include <Standard_ShortReal.hxx>
-
-#include <OpenGl_AABB.hxx>
-
-
-OpenGl_AABB::OpenGl_AABB() : myMinPoint(),
- myMaxPoint(),
- myIsValid (false)
-{ }
-
-OpenGl_AABB::OpenGl_AABB (const OpenGl_RTVec4f& thePoint) : myMinPoint (thePoint),
- myMaxPoint (thePoint),
- myIsValid (true)
-{ }
-
-OpenGl_AABB::OpenGl_AABB (const OpenGl_RTVec4f& theMinPoint,
- const OpenGl_RTVec4f& theMaxPoint) : myMinPoint (theMinPoint),
- myMaxPoint (theMaxPoint),
- myIsValid (true)
-{ }
-
-OpenGl_AABB::OpenGl_AABB (const OpenGl_AABB& theVolume) : myMinPoint (theVolume.myMinPoint),
- myMaxPoint (theVolume.myMaxPoint),
- myIsValid (theVolume.myIsValid)
-{ }
-
-void OpenGl_AABB::Add (const OpenGl_RTVec4f& thePoint)
-{
- if (!myIsValid)
- {
- myMinPoint = thePoint;
- myMaxPoint = thePoint;
- myIsValid = true;
- }
- else
- {
- myMinPoint = OpenGl_RTVec4f (Min (myMinPoint.x(), thePoint.x()),
- Min (myMinPoint.y(), thePoint.y()),
- Min (myMinPoint.z(), thePoint.z()),
- 1.f);
-
- myMaxPoint = OpenGl_RTVec4f (Max (myMaxPoint.x(), thePoint.x()),
- Max (myMaxPoint.y(), thePoint.y()),
- Max (myMaxPoint.z(), thePoint.z()),
- 1.f);
- }
-}
-
-void OpenGl_AABB::Combine (const OpenGl_AABB& theVolume)
-{
- if (!theVolume.myIsValid)
- return;
-
- if (!myIsValid)
- {
- myMinPoint = theVolume.myMinPoint;
- myMaxPoint = theVolume.myMaxPoint;
- myIsValid = true;
- }
- else
- {
- myMinPoint = OpenGl_RTVec4f (Min (myMinPoint.x(), theVolume.myMinPoint.x()),
- Min (myMinPoint.y(), theVolume.myMinPoint.y()),
- Min (myMinPoint.z(), theVolume.myMinPoint.z()),
- 1.f);
-
- myMaxPoint = OpenGl_RTVec4f (Max (myMaxPoint.x(), theVolume.myMaxPoint.x()),
- Max (myMaxPoint.y(), theVolume.myMaxPoint.y()),
- Max (myMaxPoint.z(), theVolume.myMaxPoint.z()),
- 1.f);
- }
-}
-
-OpenGl_AABB OpenGl_AABB::Added (const OpenGl_RTVec4f& thePoint) const
-{
- OpenGl_AABB result (*this);
-
- result.Add (thePoint);
-
- return result;
-}
-
-OpenGl_AABB OpenGl_AABB::Combined (const OpenGl_AABB& theVolume) const
-{
- OpenGl_AABB result (*this);
-
- result.Combine (theVolume);
-
- return result;
-}
-
-void OpenGl_AABB::Clear()
-{
- myIsValid = false;
-}
-
-OpenGl_RTVec4f OpenGl_AABB::Size() const
-{
- return myMaxPoint - myMinPoint;
-}
-
-float OpenGl_AABB::Area() const
-{
- const float aXLen = myMaxPoint.x() - myMinPoint.x();
- const float aYLen = myMaxPoint.y() - myMinPoint.y();
- const float aZLen = myMaxPoint.z() - myMinPoint.z();
-
- return ( aXLen * aYLen + aXLen * aZLen + aZLen * aYLen ) * 2.f;
-}
-
-#endif
+++ /dev/null
-// Created on: 2013-08-27
-// Created by: Denis BOGOLEPOV
-// Copyright (c) 2013 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 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_AABB_Header
-#define _OpenGl_AABB_Header
-
-#include <OpenGl_RaytraceTypes.hxx>
-
-
-//! Axis aligned bounding box (AABB).
-class OpenGl_AABB
-{
-public:
-
- //! Creates default (invalid) bounding volume.
- OpenGl_AABB();
-
- //! Creates bounding volume of given point.
- OpenGl_AABB (const OpenGl_RTVec4f& thePoint);
-
- //! Creates copy of another bounding volume.
- OpenGl_AABB (const OpenGl_AABB& theVolume);
-
- //! Creates bounding volume from min and max points.
- OpenGl_AABB (const OpenGl_RTVec4f& theMinPoint,
- const OpenGl_RTVec4f& theMaxPoint);
-
- //! Is object represents uninitialized volume?
- bool IsVoid() const { return !myIsValid; }
-
- //! Appends new point to the volume.
- void Add (const OpenGl_RTVec4f& theVector);
- //! Combines the volume with another volume.
- void Combine (const OpenGl_AABB& theVolume);
-
- //! Returns new volume created by appending a point to current volume.
- OpenGl_AABB Added (const OpenGl_RTVec4f& thePoint) const;
- //! Returns new volume created by combining with specified volume.
- OpenGl_AABB Combined (const OpenGl_AABB& theVolume) const;
-
- //! Clears bounding volume (makes object invalid).
- void Clear();
-
- //! Evaluates surface area of bounding volume.
- float Area() const;
-
- //! Return diagonal of bounding volume.
- OpenGl_RTVec4f Size() const;
-
- //! Returns minimum point of bounding volume.
- const OpenGl_RTVec4f& CornerMin() const { return myMinPoint; }
- //! Returns maximum point of bounding volume.
- const OpenGl_RTVec4f& CornerMax() const { return myMaxPoint; }
-
- //! Returns minimum point of bounding volume.
- OpenGl_RTVec4f& CornerMin() { return myMinPoint; }
- //! Returns maximum point of bounding volume.
- OpenGl_RTVec4f& CornerMax() { return myMaxPoint; }
-
-private:
-
- //! Minimum point of bounding volume.
- OpenGl_RTVec4f myMinPoint;
- //! Maximum point of bounding volume.
- OpenGl_RTVec4f myMaxPoint;
-
- //! Is bounding volume valid (up to date)?
- bool myIsValid;
-};
-
-#endif
myReleaseQueue->Pop();
}
- // release delayed shared resoruces
+ // release delayed shared resources
NCollection_Vector<TCollection_AsciiString> aDeadList;
for (NCollection_DataMap<TCollection_AsciiString, Standard_Integer>::Iterator anIter (*myDelayed);
anIter.More(); anIter.Next())
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
- EOL
/////////////////////////////////////////////////////////////////////////////////////////
// Support functions
EOL
/////////////////////////////////////////////////////////////////////////////////////////
// Functions for compute ray-object intersection
EOL
- EOL" #define _OOEPS_ exp2( -80.0f )"
+ EOL" #define _OOEPS_ exp2 (-80.0f)"
EOL
// =======================================================================
// function : IntersectSphere
// purpose : Computes ray-sphere intersection
// =======================================================================
- EOL" bool IntersectSphere (const SRay* theRay, float theRadius, float* theTime)"
+ 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
EOL" if (aD > 0.f)"
EOL" {"
- EOL" *theTime = (-aDdotO + native_sqrt (aD)) * (1.f / aDdotD);"
+ EOL" float aTime = (-aDdotO + native_sqrt (aD)) * (1.f / aDdotD);"
EOL
- EOL" return *theTime > 0.f;"
+ EOL" return aTime > 0.f ? aTime : MAXFLOAT;"
EOL" }"
EOL
- EOL" return false;"
+ EOL" return MAXFLOAT;"
EOL" }"
EOL
// =======================================================================
// function : IntersectBox
// purpose : Computes ray-box intersection (slab test)
// =======================================================================
- EOL" bool IntersectBox (const SRay* theRay,"
- EOL" float4 theMinPoint,"
- EOL" float4 theMaxPoint,"
- EOL" float* theTimeStart,"
- EOL" float* theTimeFinal)"
+ 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" const float4 aTimeMax = max (aTime0, aTime1);"
EOL" const float4 aTimeMin = min (aTime0, aTime1);"
EOL
- EOL" *theTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
- EOL" *theTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
+ 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);"
+ EOL" return (theTimeStart <= theTimeFinal) && (theTimeFinal >= 0.f) ? theTimeStart : MAXFLOAT;"
EOL" }"
EOL
// =======================================================================
// function : IntersectTriangle
// purpose : Computes ray-triangle intersection (branchless version)
// =======================================================================
- EOL" bool IntersectTriangle (const SRay* theRay,"
- EOL" const float4 thePoint0,"
- EOL" const float4 thePoint1,"
- EOL" const float4 thePoint2,"
- EOL" float4* theNormal,"
- EOL" float* theTime,"
- EOL" float* theU,"
- EOL" float* theV)"
+ 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" const float4 aEdge2 = (1.f / dot (*theNormal, theRay->Direct)) * (thePoint0 - theRay->Origin);"
EOL
- EOL" *theTime = dot (*theNormal, aEdge2);"
+ EOL" const float aTime = dot (*theNormal, aEdge2);"
EOL
- EOL" const float4 theInc = cross (theRay->Direct, aEdge2);"
+ EOL" const float4 theVec = cross (theRay->Direct, aEdge2);"
EOL
- EOL" *theU = dot (theInc, aEdge1);"
- EOL" *theV = dot (theInc, aEdge0);"
+ EOL" *theU = dot (theVec, aEdge1);"
+ EOL" *theV = dot (theVec, aEdge0);"
EOL
- EOL" return (*theTime > 0) & (*theU >= 0.f) & (*theV >= 0.f) & (*theU + *theV <= 1.f);"
+ EOL" return (aTime >= 0.f) & (*theU >= 0.f) & (*theV >= 0.f) & (*theU + *theV <= 1.f) ? aTime : MAXFLOAT;"
EOL" }"
EOL
/////////////////////////////////////////////////////////////////////////////////////////
// function : Shade
// purpose : Computes Phong-based illumination
// =======================================================================
- EOL" float4 Shade (__global float4* theMaterials,"
+ 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" const int theMatIndex)"
+ EOL" const float theTranspr)"
EOL" {"
EOL" float aLambert = dot (theNormal, theLight);"
EOL
EOL
EOL" if (aLambert > 0.f)"
EOL" {"
- EOL" const float4 aMatDiff = theMaterials[7 * theMatIndex + 1];"
- EOL" const float4 aMatSpec = theMaterials[7 * theMatIndex + 2];"
- 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), aMatSpec.w);"
+ EOL" const float aSpecular = pow (max (dot (aReflect.xyz, theView.xyz), 0.f), theMatSpec.w);"
EOL
- EOL" return theIntens * (aMatDiff * aLambert + aMatSpec * aSpecular);"
+ EOL" return theIntens * (theMatDiff * aLambert + theMatSpec * aSpecular);"
EOL" }"
EOL
EOL" return ZERO;"
EOL" }"
EOL
// =======================================================================
- // function : Lat-long
+ // function : Latlong
// purpose : Converts world direction to environment texture coordinates
// =======================================================================
- EOL" float2 Latlong (const float4 theDirect)"
+ EOL" float2 Latlong (const float4 thePoint, const float theRadius)"
EOL" {"
- EOL" float aPsi = acos( -theDirect.y );"
- EOL" float aPhi = atan2( theDirect.z, theDirect.x );"
+ EOL" float aPsi = acospi (-thePoint.y / theRadius);"
+ EOL" float aPhi = atan2pi (thePoint.z, thePoint.x);"
EOL
- EOL" aPhi = (aPhi < 0) ? (aPhi + 2.f * M_PI_F) : aPhi;"
+ EOL" aPhi = (aPhi < 0.f) ? aPhi + 2.f : aPhi;"
EOL
- EOL" return (float2) (aPhi / (2.f * M_PI_F), aPsi / M_PI_F);"
+ EOL" return (float2) (aPhi * 0.5f, aPsi);"
EOL" }"
EOL
/////////////////////////////////////////////////////////////////////////////////////////
EOL" (*thePos)--;"
EOL" }"
EOL
- // #define BVH_MINIMIZE_MEM_LOADS
- EOL
// =======================================================================
- // function : Traverse
- // purpose : Finds intersection with nearest triangle
+ // function : ObjectNearestHit
+ // purpose : Finds intersection with nearest object triangle
// =======================================================================
- EOL" int4 Traverse (const SRay* theRay,"
- EOL" __global int4* theIndices,"
- EOL" __global float4* theVertices,"
- EOL" __global float4* theNodeMinPoints,"
- EOL" __global float4* theNodeMaxPoints,"
- EOL" __global int4* theNodeDataRecords,"
- EOL" SIntersect* theHit)"
+ 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" char aHead = -1;"
EOL
- EOL" uint aNode = 0;" // root node
+ 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" float4 aNodeMinLft;"
- EOL" float4 aNodeMaxLft;"
- EOL" float4 aNodeMinRgh;"
- EOL" float4 aNodeMaxRgh;"
+ EOL" while (true)"
+ EOL" {"
+ EOL" const int3 aData = theObjectNodeInfoBuffer[aNode].xyz;"
EOL
- EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
- EOL" aNodeMinLft = theNodeMinPoints[aNode];"
- EOL" aNodeMaxLft = theNodeMaxPoints[aNode];"
- EOL" #endif"
+ EOL" if (aData.x == 0)" // if inner node
+ EOL" {"
+ EOL" float4 aNodeMin = theObjectMinPointBuffer[aData.y];"
+ EOL" float4 aNodeMax = theObjectMaxPointBuffer[aData.y];"
EOL
- EOL" int4 aTriangleIndex = (int4) (-1);"
+ EOL" float4 aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
+ EOL" float4 aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
EOL
- EOL" theHit->Time = MAXFLOAT;"
+ EOL" float4 aTimeMax = max (aTime0, aTime1);"
+ EOL" float4 aTimeMin = min (aTime0, aTime1);"
EOL
- EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
- EOL" int3 aData = (int3) (1,"
- EOL" as_int (aNodeMinLft.w),"
- EOL" as_int (aNodeMaxLft.w));"
+ EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
+ EOL" aTimeMin1 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
EOL
- EOL" aData = aData.y < 0 ? -aData : aData;"
- EOL" #endif"
+ EOL" const bool aHitLft = (aTimeMin1 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin1 <= theIntersect->Time);"
EOL
- EOL" while (true)"
- EOL" {"
- EOL" #ifndef BVH_MINIMIZE_MEM_LOADS"
- EOL" int3 aData = theNodeDataRecords[aNode].xyz;"
- EOL" #endif"
+ EOL" aNodeMin = theObjectMinPointBuffer[aData.z];"
+ EOL" aNodeMax = theObjectMaxPointBuffer[aData.z];"
EOL
- EOL" if (aData.x != 1)" // if inner node
- EOL" {"
- EOL" aNodeMinLft = theNodeMinPoints[aData.y];"
- EOL" aNodeMinRgh = theNodeMinPoints[aData.z];"
- EOL" aNodeMaxLft = theNodeMaxPoints[aData.y];"
- EOL" aNodeMaxRgh = theNodeMaxPoints[aData.z];"
+ EOL" aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
+ EOL" aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
EOL
- EOL" IntersectNodes (theRay,"
- EOL" aNodeMinLft,"
- EOL" aNodeMaxLft,"
- EOL" aNodeMinRgh,"
- EOL" aNodeMaxRgh,"
- EOL" &aTimeMin1,"
- EOL" &aTimeMin2,"
- EOL" theHit->Time);"
+ EOL" aTimeMax = max (aTime0, aTime1);"
+ EOL" aTimeMin = min (aTime0, aTime1);"
EOL
- EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
- EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
+ 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" #ifdef BVH_MINIMIZE_MEM_LOADS"
- EOL" aData = (int3) (1,"
- EOL" as_int (aTimeMin1 < aTimeMin2 ? aNodeMinLft.w : aNodeMinRgh.w),"
- EOL" as_int (aTimeMin1 < aTimeMin2 ? aNodeMaxLft.w : aNodeMaxRgh.w));"
- EOL
- EOL" aData = aData.y < 0 ? -aData : aData;"
- EOL" #endif"
EOL" }"
EOL" else"
EOL" {"
EOL" if (aHitLft | aHitRgh)"
EOL" {"
EOL" aNode = aHitLft ? aData.y : aData.z;"
- EOL
- EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
- EOL" aData = (int3) (1,"
- EOL" as_int (aHitLft ? aNodeMinLft.w : aNodeMinRgh.w),"
- EOL" as_int (aHitLft ? aNodeMaxLft.w : aNodeMaxRgh.w));"
- EOL
- EOL" aData = aData.y < 0 ? -aData : aData;"
- EOL" #endif"
EOL" }"
EOL" else"
EOL" {"
EOL" return aTriangleIndex;"
EOL
EOL" pop (aStack, &aHead, &aNode);"
- EOL
- EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
- EOL" aData = theNodeDataRecords[aNode].xyz;"
- EOL" #endif"
EOL" }"
EOL" }"
EOL" }"
EOL" else " // if leaf node
EOL" {"
- EOL" for (int nTri = aData.y; nTri <= aData.z; ++nTri)"
+ EOL" for (int aTriIdx = aData.y; aTriIdx <= aData.z; ++aTriIdx)"
EOL" {"
- EOL" int4 anIndex = theIndices[nTri];"
+ EOL" const int4 aTestTriangle = theGeometryTriangBuffer[aTriIdx];"
EOL
- EOL" const float4 aP0 = theVertices[anIndex.x];"
- EOL" const float4 aP1 = theVertices[anIndex.y];"
- EOL" const float4 aP2 = theVertices[anIndex.z];"
+ 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;"
+ EOL" float4 aNormal; float aU, aV;"
EOL
- EOL" float aTime, aU, aV;"
+ EOL" float aTime = IntersectTriangle (theRay,"
+ EOL" aPoint0,"
+ EOL" aPoint1,"
+ EOL" aPoint2,"
+ EOL" &aNormal,"
+ EOL" &aU,"
+ EOL" &aV);"
EOL
- EOL" if (IntersectTriangle (theRay, aP0, aP1, aP2, &aNormal, &aTime, &aU, &aV) & (aTime < theHit->Time))"
+ EOL" if (aTime < theIntersect->Time)"
EOL" {"
- EOL" aTriangleIndex = anIndex;"
- EOL" theHit->Normal = aNormal;"
- EOL" theHit->Time = aTime;"
- EOL" theHit->U = aU;"
- EOL" theHit->V = aV;"
+ EOL" aTriangleIndex = aTestTriangle;"
+ EOL" theIntersect->Normal = aNormal;"
+ EOL" theIntersect->Time = aTime;"
+ EOL" theIntersect->U = aU;"
+ EOL" theIntersect->V = aV;"
EOL" }"
EOL" }"
EOL
EOL" return aTriangleIndex;"
EOL
EOL" pop (aStack, &aHead, &aNode);"
- EOL
- EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
- EOL" aData = theNodeDataRecords[aNode].xyz;"
- EOL" #endif"
EOL" }"
EOL" }"
EOL
EOL" return aTriangleIndex;"
- EOL" }"
- EOL
- EOL" #define TRANSPARENT_SHADOW_"
+ EOL" }"
EOL
// =======================================================================
- // function : TraverseShadow
- // purpose : Finds intersection with any triangle
+ // function : ObjectAnyHit
+ // purpose : Finds intersection with any object triangle
// =======================================================================
- EOL" float TraverseShadow (const SRay* theRay,"
- EOL" __global int4* theIndices,"
- EOL" __global float4* theVertices,"
- EOL" __global float4* materials,"
- EOL" __global float4* theNodeMinPoints,"
- EOL" __global float4* theNodeMaxPoints,"
- EOL" __global int4* theNodeDataRecords,"
- EOL" float theDistance)"
+ 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" char aHead = -1;"
EOL
- EOL" uint aNode = 0;" // root node
+ EOL" char aHead = -1;" // stack pointer
+ EOL" uint aNode = 0;" // node to visit
EOL
- EOL" float aFactor = 1.f;" // light attenuation factor
+ 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" int3 aData = theNodeDataRecords[aNode].xyz;"
+ 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 != 1)" // if inner node
+ 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" theNodeMinPoints[aData.y],"
- EOL" theNodeMaxPoints[aData.y],"
- EOL" theNodeMinPoints[aData.z],"
- EOL" theNodeMaxPoints[aData.z],"
+ EOL" aNodeMinLft,"
+ EOL" aNodeMaxLft,"
+ EOL" aNodeMinRgh,"
+ EOL" aNodeMaxRgh,"
EOL" &aTimeMin1,"
EOL" &aTimeMin2,"
- EOL" theDistance);"
+ EOL" theIntersect->Time);"
EOL
EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
EOL" else"
EOL" {"
EOL" if (aHead < 0)"
- EOL" return aFactor;"
+ EOL" return aNearestTriangle;"
EOL
EOL" pop (aStack, &aHead, &aNode);"
EOL" }"
EOL" }"
EOL" }"
- EOL" else " // if leaf node
- EOL" {"
- EOL" for (int nTri = aData.y; nTri <= aData.z; ++nTri)"
- EOL" {"
- EOL" int4 anIndex = theIndices[nTri];"
+ 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" const float4 aP0 = theVertices[anIndex.x];"
- EOL" const float4 aP1 = theVertices[anIndex.y];"
- EOL" const float4 aP2 = theVertices[anIndex.z];"
+ EOL" char aHead = -1;" // stack pointer
+ EOL" uint aNode = 0;" // node to visit
EOL
- EOL" float4 aNormal;"
+ EOL" while (true)"
+ EOL" {"
+ EOL" const int4 aData = theSceneNodeInfoBuffer[aNode];"
EOL
- EOL" float aTime, aU, aV;"
+ EOL" if (aData.x != 0)" // if leaf node
+ EOL" {"
+ EOL" const float4 aNodeMin = theSceneMinPointBuffer[aNode];"
+ EOL" const float4 aNodeMax = theSceneMaxPointBuffer[aNode];"
EOL
- EOL" if (IntersectTriangle (theRay, aP0, aP1, aP2, &aNormal, &aTime, &aU, &aV) & (aTime < theDistance))"
+ 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" #ifdef TRANSPARENT_SHADOW"
- EOL" aFactor *= materials[7 * index.w + 6].x;"
- EOL
- EOL" if (aFactor < 0.1f)"
- EOL" return aFactor;"
- EOL" #else"
EOL" return 0.f;"
- EOL" #endif"
EOL" }"
EOL" }"
EOL
EOL" if (aHead < 0)"
- EOL" return aFactor;"
+ EOL" return 1.f;"
EOL
EOL" pop (aStack, &aHead, &aNode);"
EOL" }"
- 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" return aFactor;"
+ 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 _MAT_SIZE_ 7"
+ EOL" #define THRESHOLD (float4) (0.1f, 0.1f, 0.1f, 1.f)"
EOL
- EOL" #define _LGH_SIZE_ 3"
+ 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 : Raytrace
+ // function : Radiance
// purpose : Computes color of specified ray
// =======================================================================
- EOL" float4 Raytrace (SRay* theRay,"
+ EOL" float4 Radiance (SRay* theRay,"
EOL" __read_only image2d_t theEnvMap,"
- EOL" __global float4* theNodeMinPoints,"
- EOL" __global float4* theNodeMaxPoints,"
- EOL" __global int4* theNodeDataRecords,"
- EOL" __global float4* theLightSources,"
- EOL" __global float4* theMaterials,"
- EOL" __global float4* theVertices,"
- EOL" __global float4* theNormals,"
- EOL" __global int4* theIndices,"
- EOL" const int theLightCount,"
- EOL" const float theEpsilon,"
- EOL" const float theRadius,"
- EOL" const int isShadows,"
- EOL" const int isReflect)"
+ 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" for (int aDepth = 0; aDepth < _MAX_DEPTH_; ++aDepth)"
EOL" {"
- EOL" int4 aTriangle = Traverse (theRay,"
- EOL" theIndices,"
- EOL" theVertices,"
- EOL" theNodeMinPoints,"
- EOL" theNodeMaxPoints,"
- EOL" theNodeDataRecords,"
- EOL" &aHit);"
+ 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" float aTime;"
- EOL
- EOL" if (aWeight.w != 0.f || !IntersectSphere (theRay, theRadius, &aTime))"
+ EOL" if (aWeight.w != 0.f)"
EOL" break;"
EOL
- EOL" float2 aTexCoord = Latlong (fma (theRay->Direct, (float4) (aTime), theRay->Origin) * (1.f / theRadius));"
+ EOL" float aTime = IntersectSphere (theRay, theSceneRadius);"
EOL
- EOL" aResult += aWeight * read_imagef (theEnvMap, EnvironmentSampler, aTexCoord);"
+ 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" aWeight.w);"
EOL" }"
EOL
- EOL" " // Compute geometric normal
+ EOL // Compute geometric normal
EOL" float4 aGeomNormal = aHit.Normal; aGeomNormal = fast_normalize (aGeomNormal);"
EOL
- EOL" " // Compute interpolated normal
- EOL" float4 aNormal = SmoothNormal (theNormals, &aHit, aTriangle);"
+ EOL // Compute interpolated normal
+ EOL" float4 aNormal = SmoothNormal (theGeometryNormalBuffer, &aHit, aTriangle);"
EOL
- EOL" " // Compute intersection point
+ EOL // Compute intersection point
EOL" float4 aPoint = theRay->Direct * aHit.Time + theRay->Origin;"
EOL
- EOL" float4 aMaterAmb = theMaterials [_MAT_SIZE_ * aTriangle.w + 0];"
- EOL" float4 aMaterTrn = theMaterials [_MAT_SIZE_ * aTriangle.w + 6];"
- EOL
- EOL" for (int nLight = 0; nLight < theLightCount; ++nLight)"
- EOL" {"
- EOL" float4 aLightAmbient = theLightSources [_LGH_SIZE_ * nLight];"
+ EOL" float4 aMaterAmb = MATERIAL_AMBN (theMaterialBuffer, aTriangle);"
+ EOL" float4 aMaterTrn = MATERIAL_TRAN (theMaterialBuffer, aTriangle);"
EOL
- EOL" aResult += aWeight * aLightAmbient * aMaterAmb *"
+ EOL" aResult += aWeight * theGlobalAmbient * aMaterAmb *"
EOL" (aMaterTrn.x * max (fabs (dot (theRay->Direct, aNormal)), 0.5f));"
EOL
- EOL" if (aLightAmbient.w < 0.f)" // 'ambient' light
- EOL" {"
- EOL" continue;" // 'ambient' light has no another luminances
- EOL" }"
- EOL
- EOL" float4 aLightPosition = theLightSources [_LGH_SIZE_ * nLight + 2];"
+ 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" aShadow.Direct = (aLightPosition - aPoint) * (1.f / aLightDistance);"
EOL" }"
EOL
- EOL" aShadow.Origin = aPoint + aShadow.Direct * theEpsilon +"
- EOL" aGeomNormal * copysign (theEpsilon, dot (aGeomNormal, aShadow.Direct));"
+ EOL" aShadow.Origin = aPoint + aShadow.Direct * theSceneEpsilon +"
+ EOL" aGeomNormal * copysign (theSceneEpsilon, dot (aGeomNormal, aShadow.Direct));"
EOL
- EOL" float aFactor = 1.f;"
+ EOL" float aVisibility = 1.f;"
EOL
- EOL" if (isShadows)"
+ EOL" if (theShadowsEnabled)"
EOL" {"
- EOL" aFactor = TraverseShadow (&aShadow,"
- EOL" theIndices,"
- EOL" theVertices,"
- EOL" theMaterials,"
- EOL" theNodeMinPoints,"
- EOL" theNodeMaxPoints,"
- EOL" theNodeDataRecords,"
- EOL" aLightDistance);"
+ 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" aResult += (aMaterTrn.x * aFactor) * aWeight * Shade (theMaterials,"
- EOL" aShadow.Direct,"
- EOL" -theRay->Direct,"
- EOL" aNormal,"
- EOL" UNIT,"
- EOL" aMaterTrn.y,"
- EOL" aTriangle.w);"
+ 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" else"
EOL" {"
- EOL" float4 aMaterRef = theMaterials [_MAT_SIZE_ * aTriangle.w + 4];"
- EOL" aWeight *= isReflect ? aMaterRef : ZERO;"
+ EOL" aWeight *= theReflectEnabled ? MATERIAL_REFL (theMaterialBuffer, aTriangle) : ZERO;"
EOL
- EOL" theRay->Direct -= 2.f * dot (theRay->Direct, aNormal) * aNormal;"
+ EOL" float4 aDirect = theRay->Direct - 2.f * dot (theRay->Direct, aNormal) * aNormal;"
EOL
- EOL" float aDdotN = dot (theRay->Direct, aGeomNormal);"
+ EOL" float aDdotN = dot (aDirect, aGeomNormal);"
EOL" if (aDdotN < 0.f)"
- EOL" theRay->Direct -= aDdotN * aGeomNormal;"
+ EOL" theRay->Direct -= 2.f * dot (theRay->Direct, aGeomNormal) * aGeomNormal;"
+ EOL" else"
+ EOL" theRay->Direct = aDirect;"
EOL" }"
EOL
- EOL" if (aWeight.x < 0.1f && aWeight.y < 0.1f && aWeight.z < 0.1f)"
+ EOL" if (all (islessequal (aWeight, THRESHOLD)))"
EOL" {"
EOL" return (float4) (aResult.x,"
EOL" aResult.y,"
EOL" aWeight.w);"
EOL" }"
EOL
- EOL" theRay->Origin = theRay->Direct * theEpsilon + aPoint;"
+ EOL" theRay->Origin = theRay->Direct * theSceneEpsilon + aPoint;"
EOL" }"
EOL
EOL" return (float4) (aResult.x,"
EOL" aWeight.w);"
EOL" }"
EOL
- EOL
///////////////////////////////////////////////////////////////////////////////
// Ray tracing kernel functions
EOL
// =======================================================================
- // function : Main
+ // function : RaytraceRender
// purpose : Computes pixel color using ray-tracing
// =======================================================================
- EOL" __kernel void Main (__write_only image2d_t theOutput,"
- EOL" __read_only image2d_t theEnvMap,"
- EOL" __global float4* theNodeMinPoints,"
- EOL" __global float4* theNodeMaxPoints,"
- EOL" __global int4* theNodeDataRecords,"
- EOL" __global float4* theLightSources,"
- EOL" __global float4* theMaterials,"
- EOL" __global float4* theVertices,"
- EOL" __global float4* theNormals,"
- EOL" __global int4* theIndices,"
- EOL" const float16 theOrigins,"
- EOL" const float16 theDirects,"
- EOL" const int theLightCount,"
- EOL" const float theEpsilon,"
- EOL" const float theRadius,"
- EOL" const int isShadows,"
- EOL" const int isReflect,"
- EOL" const int theSizeX,"
- EOL" const int theSizeY)"
+ 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 aX = get_global_id (0);"
- EOL" const int aY = get_global_id (1);"
+ EOL" const int aPixelX = get_global_id (0);"
+ EOL" const int aPixelY = get_global_id (1);"
EOL
- EOL" if (aX >= theSizeX || aY >= theSizeY)"
+ EOL" if (aPixelX >= theSizeX || aPixelY >= theSizeY)"
EOL" return;"
EOL
EOL" private SRay aRay;"
EOL
EOL" GenerateRay (&aRay,"
- EOL" aX,"
- EOL" aY,"
+ EOL" aPixelX,"
+ EOL" aPixelY,"
EOL" theSizeX,"
EOL" theSizeY,"
EOL" theOrigins,"
EOL" theDirects);"
EOL
- EOL" float4 aBoxMin = theNodeMinPoints[0] - (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
- EOL" float4 aBoxMax = theNodeMaxPoints[0] + (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
- EOL
- EOL" float aTimeStart;"
- EOL" float aTimeFinal;"
- EOL
EOL" float4 aColor = (float4) (0.f, 0.f, 0.f, 1.f);"
EOL
- EOL" if (IntersectBox (&aRay, aBoxMin, aBoxMax, &aTimeStart, &aTimeFinal))"
+ EOL" float aTimeStart = IntersectBox (&aRay, theSceneMinPointBuffer[0], theSceneMaxPointBuffer[0]);"
+ EOL
+ EOL" if (aTimeStart != MAXFLOAT)"
EOL" {"
- EOL" aRay.Origin = fma (aRay.Direct, (float4) (aTimeStart), aRay.Origin);"
- EOL
- EOL" aColor = Raytrace (&aRay,"
- EOL" theEnvMap,"
- EOL" theNodeMinPoints,"
- EOL" theNodeMaxPoints,"
- EOL" theNodeDataRecords,"
- EOL" theLightSources,"
- EOL" theMaterials,"
- EOL" theVertices,"
- EOL" theNormals,"
- EOL" theIndices,"
- EOL" theLightCount,"
- EOL" theEpsilon,"
- EOL" theRadius,"
- EOL" isShadows,"
- EOL" isReflect);"
+ 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) (aX, aY), aColor);"
+ 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.075f"
+ EOL" #define _LUM_DELTA_ 0.085f"
EOL
EOL" #define AA_MAX 0.559017f"
EOL" #define AA_MIN 0.186339f"
EOL
// =======================================================================
- // function : MainAntialiased
+ // function : RaytraceSmooth
// purpose : Performs adaptive sub-pixel rendering
// =======================================================================
- EOL" __kernel void MainAntialiased ( __read_only image2d_t theInput,"
- EOL" __write_only image2d_t theOutput,"
- EOL" __read_only image2d_t theEnvMap,"
- EOL" __global float4* theNodeMinPoints,"
- EOL" __global float4* theNodeMaxPoints,"
- EOL" __global int4* theNodeDataRecords,"
- EOL" __global float4* theLightSources,"
- EOL" __global float4* theMaterials,"
- EOL" __global float4* theVertices,"
- EOL" __global float4* theNormals,"
- EOL" __global int4* theIndices,"
- EOL" const float16 theOrigins,"
- EOL" const float16 theDirects,"
- EOL" const int theLightCount,"
- EOL" const float theEpsilon,"
- EOL" const float theRadius,"
- EOL" const int isShadows,"
- EOL" const int isReflect,"
- EOL" const int theSizeX,"
- EOL" const int theSizeY )"
+ 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 aX = get_global_id (0);"
- EOL" const int aY = get_global_id (1);"
+ EOL" const int aPixelX = get_global_id (0);"
+ EOL" const int aPixelY = get_global_id (1);"
EOL
- EOL" if (aX >= theSizeX || aY >= theSizeY)"
+ EOL" if (aPixelX >= theSizeX || aPixelY >= theSizeY)"
EOL" return;"
EOL
- EOL" float4 aClr0 = read_imagef (theInput, OutputSampler, (float2) (aX + 0, aY + 0));"
- EOL" float4 aClr1 = read_imagef (theInput, OutputSampler, (float2) (aX + 0, aY - 1));"
- EOL" float4 aClr2 = read_imagef (theInput, OutputSampler, (float2) (aX + 0, aY + 1));"
- EOL
- EOL" float4 aClr3 = read_imagef (theInput, OutputSampler, (float2) (aX + 1, aY + 0));"
- EOL" float4 aClr4 = read_imagef (theInput, OutputSampler, (float2) (aX + 1, aY - 1));"
- EOL" float4 aClr5 = read_imagef (theInput, OutputSampler, (float2) (aX + 1, aY + 1));"
+ 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 aClr6 = read_imagef (theInput, OutputSampler, (float2) (aX - 1, aY + 0));"
- EOL" float4 aClr7 = read_imagef (theInput, OutputSampler, (float2) (aX - 1, aY - 1));"
- EOL" float4 aClr8 = read_imagef (theInput, OutputSampler, (float2) (aX - 1, aY + 1));"
+ 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" 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" 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" float aLum = (aClr0.w == 1.f) ? -1.f : (0.2126f * aClr0.x + 0.7152f * aClr0.y + 0.0722f * aClr0.z);"
+ 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" bool 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" float4 aColor = aClr0;"
+ EOL" float4 aColor = clamp (aClr0, 0.f, 1.f);"
EOL
EOL" private SRay aRay;"
EOL
- EOL" const float4 aBoxMin = theNodeMinPoints[0] - (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
- EOL" const float4 aBoxMax = theNodeMaxPoints[0] + (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
+ 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 = aX, fY = aY;"
+ EOL" float fX = aPixelX, fY = aPixelY;"
EOL
EOL" if (aSample == 0)"
EOL" {"
EOL" theOrigins,"
EOL" theDirects);"
EOL
- EOL" float aTimeStart;"
- EOL" float aTimeFinal;"
+ EOL" float aTimeStart = IntersectBox (&aRay, aBoxMin, aBoxMax);"
EOL
- EOL" if (IntersectBox (&aRay, aBoxMin, aBoxMax, &aTimeStart, &aTimeFinal))"
+ EOL" if (aTimeStart != MAXFLOAT)"
EOL" {"
- EOL" aRay.Origin = fma (aRay.Direct, (float4) (aTimeStart), aRay.Origin);"
- EOL
- EOL" aColor += Raytrace (&aRay,"
- EOL" theEnvMap,"
- EOL" theNodeMinPoints,"
- EOL" theNodeMaxPoints,"
- EOL" theNodeDataRecords,"
- EOL" theLightSources,"
- EOL" theMaterials,"
- EOL" theVertices,"
- EOL" theNormals,"
- EOL" theIndices,"
- EOL" theLightCount,"
- EOL" theEpsilon,"
- EOL" theRadius,"
- EOL" isShadows,"
- EOL" isReflect);"
+ 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" aColor *= 1.f / 5.f;"
EOL" }"
EOL
- EOL" write_imagef (theOutput, (int2) (aX, aY), aColor);"
+ EOL" write_imagef (theOutput, (int2) (aPixelX, aPixelY), aColor);"
EOL" }";
#endif
+++ /dev/null
-// Created on: 2013-10-15
-// Created by: Denis BOGOLEPOV
-// Copyright (c) 2013 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 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_RaytraceTypes_Header
-#define _OpenGl_RaytraceTypes_Header
-
-#include <vector>
-
-#include <NCollection_Vec4.hxx>
-#include <NCollection_StdAllocator.hxx>
-
-//! 4D vector of integers.
-typedef NCollection_Vec4<int> OpenGl_RTVec4i;
-
-//! 4D vector of floats.
-typedef NCollection_Vec4<float> OpenGl_RTVec4f;
-
-//! 4D vector of doubles.
-typedef NCollection_Vec4<double> OpenGl_RTVec4d;
-
-//! Array of 4D integer vectors.
-typedef std::vector<OpenGl_RTVec4i,
- NCollection_StdAllocator<OpenGl_RTVec4i> > OpenGl_RTArray4i;
-
-//! Array of 4D floating point vectors.
-typedef std::vector<OpenGl_RTVec4f,
- NCollection_StdAllocator<OpenGl_RTVec4f> > OpenGl_RTArray4f;
-
-#endif
#ifdef HAVE_OPENCL
-#include <limits>
+#include <Standard_Assert.hxx>
-#include <OpenGl_SceneGeometry.hxx>
-
-namespace
-{
+#ifdef HAVE_TBB
+ #include <tbb/tbb.h>
+#endif
- //! Number of node bins per axis.
- static const int THE_NUMBER_OF_BINS = 32;
+#include <OpenGl_SceneGeometry.hxx>
- //! Max number of triangles per leaf node.
- static const int THE_MAX_LEAF_TRIANGLES = 4;
+//! Use this macro to output BVH profiling info
+//#define BVH_PRINT_INFO
- //! Useful constant for null integer 4D vector.
- static const OpenGl_RTVec4i THE_ZERO_VEC_4I;
+#ifdef BVH_PRINT_INFO
+ #include <OSD_Timer.hxx>
+#endif
+namespace
+{
//! Useful constant for null floating-point 4D vector.
- static const OpenGl_RTVec4f THE_ZERO_VEC_4F;
-
+ static const BVH_Vec4f ZERO_VEC_4F;
};
// =======================================================================
// purpose : Creates new default material
// =======================================================================
OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial()
-: Ambient (THE_ZERO_VEC_4F),
- Diffuse (THE_ZERO_VEC_4F),
- Specular (THE_ZERO_VEC_4F),
- Emission (THE_ZERO_VEC_4F),
- Reflection (THE_ZERO_VEC_4F),
- Refraction (THE_ZERO_VEC_4F),
- Transparency (THE_ZERO_VEC_4F)
+: Ambient (ZERO_VEC_4F),
+ Diffuse (ZERO_VEC_4F),
+ Specular (ZERO_VEC_4F),
+ Emission (ZERO_VEC_4F),
+ Reflection (ZERO_VEC_4F),
+ Refraction (ZERO_VEC_4F),
+ Transparency (ZERO_VEC_4F)
{ }
// =======================================================================
// function : OpenGl_RaytraceMaterial
// purpose : Creates new material with specified properties
// =======================================================================
-OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient,
- const OpenGl_RTVec4f& theDiffuse,
- const OpenGl_RTVec4f& theSpecular)
+OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial (const BVH_Vec4f& theAmbient,
+ const BVH_Vec4f& theDiffuse,
+ const BVH_Vec4f& theSpecular)
: Ambient (theAmbient),
Diffuse (theDiffuse),
Specular (theSpecular),
- Emission (THE_ZERO_VEC_4F),
- Reflection (THE_ZERO_VEC_4F),
- Refraction (THE_ZERO_VEC_4F),
- Transparency (THE_ZERO_VEC_4F)
+ Emission (ZERO_VEC_4F),
+ Reflection (ZERO_VEC_4F),
+ Refraction (ZERO_VEC_4F),
+ Transparency (ZERO_VEC_4F)
{
//
}
// function : OpenGl_RaytraceMaterial
// purpose : Creates new material with specified properties
// =======================================================================
-OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient,
- const OpenGl_RTVec4f& theDiffuse,
- const OpenGl_RTVec4f& theSpecular,
- const OpenGl_RTVec4f& theEmission,
- const OpenGl_RTVec4f& theTranspar)
+OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial (const BVH_Vec4f& theAmbient,
+ const BVH_Vec4f& theDiffuse,
+ const BVH_Vec4f& theSpecular,
+ const BVH_Vec4f& theEmission,
+ const BVH_Vec4f& theTranspar)
: Ambient (theAmbient),
Diffuse (theDiffuse),
Specular (theSpecular),
Emission (theEmission),
- Reflection (THE_ZERO_VEC_4F),
- Refraction (THE_ZERO_VEC_4F),
+ Reflection (ZERO_VEC_4F),
+ Refraction (ZERO_VEC_4F),
Transparency (theTranspar)
{
//
// function : OpenGl_RaytraceMaterial
// purpose : Creates new material with specified properties
// =======================================================================
-OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient,
- const OpenGl_RTVec4f& theDiffuse,
- const OpenGl_RTVec4f& theSpecular,
- const OpenGl_RTVec4f& theEmission,
- const OpenGl_RTVec4f& theTranspar,
- const OpenGl_RTVec4f& theReflection,
- const OpenGl_RTVec4f& theRefraction)
+OpenGl_RaytraceMaterial::OpenGl_RaytraceMaterial (const BVH_Vec4f& theAmbient,
+ const BVH_Vec4f& theDiffuse,
+ const BVH_Vec4f& theSpecular,
+ const BVH_Vec4f& theEmission,
+ const BVH_Vec4f& theTranspar,
+ const BVH_Vec4f& theReflection,
+ const BVH_Vec4f& theRefraction)
: Ambient (theAmbient),
Diffuse (theDiffuse),
Specular (theSpecular),
// function : OpenGl_LightSource
// purpose : Creates new light source
// =======================================================================
-OpenGl_RaytraceLight::OpenGl_RaytraceLight (const OpenGl_RTVec4f& theAmbient)
-: Ambient (theAmbient)
-{
- //
-}
-
-// =======================================================================
-// function : OpenGl_LightSource
-// purpose : Creates new light source
-// =======================================================================
-OpenGl_RaytraceLight::OpenGl_RaytraceLight (const OpenGl_RTVec4f& theDiffuse,
- const OpenGl_RTVec4f& thePosition)
+OpenGl_RaytraceLight::OpenGl_RaytraceLight (const BVH_Vec4f& theDiffuse,
+ const BVH_Vec4f& thePosition)
: Diffuse (theDiffuse),
Position (thePosition)
{
//
}
-// =======================================================================
-// function : Center
-// purpose : Returns centroid of specified triangle
-// =======================================================================
-OpenGl_RTVec4f OpenGl_RaytraceScene::Center (const int theTriangle) const
-{
- const OpenGl_RTVec4i& anIndex = Triangles [theTriangle];
-
- return ( Vertices[anIndex.x()] +
- Vertices[anIndex.y()] +
- Vertices[anIndex.z()] ) * ( 1.f / 3.f );
-}
-
-// =======================================================================
-// function : CenterAxis
-// purpose : Returns centroid of specified triangle
-// =======================================================================
-float OpenGl_RaytraceScene::CenterAxis (const int theTriangle,
- const int theAxis) const
-{
- const OpenGl_RTVec4i& anIndex = Triangles [theTriangle];
-
- return ( Vertices[anIndex.x()][theAxis] +
- Vertices[anIndex.y()][theAxis] +
- Vertices[anIndex.z()][theAxis] ) * ( 1.f / 3.f );
-}
-
-// =======================================================================
-// function : Box
-// purpose : Returns AABB of specified triangle
-// =======================================================================
-OpenGl_AABB OpenGl_RaytraceScene::Box (const int theTriangle) const
-{
- const OpenGl_RTVec4i& anIndex = Triangles[theTriangle];
-
- const OpenGl_RTVec4f& pA = Vertices[anIndex.x()];
- const OpenGl_RTVec4f& pB = Vertices[anIndex.y()];
- const OpenGl_RTVec4f& pC = Vertices[anIndex.z()];
-
- OpenGl_AABB aBox (pA);
-
- aBox.Add (pB);
- aBox.Add (pC);
-
- return aBox;
-}
-
// =======================================================================
// function : Clear
-// purpose : Clears all scene geometry data
+// purpose : Clears ray-tracing geometry
// =======================================================================
-void OpenGl_RaytraceScene::Clear()
+void OpenGl_RaytraceGeometry::Clear()
{
- AABB.Clear();
+ BVH_Geometry<Standard_ShortReal, 4>::BVH_Geometry::Clear();
- OpenGl_RTArray4f anEmptyNormals;
- Normals.swap (anEmptyNormals);
+ std::vector<OpenGl_RaytraceLight,
+ NCollection_StdAllocator<OpenGl_RaytraceLight> > anEmptySources;
- OpenGl_RTArray4f anEmptyVertices;
- Vertices.swap (anEmptyVertices);
-
- OpenGl_RTArray4i anEmptyTriangles;
- Triangles.swap (anEmptyTriangles);
+ Sources.swap (anEmptySources);
std::vector<OpenGl_RaytraceMaterial,
- NCollection_StdAllocator<OpenGl_RaytraceMaterial> > anEmptyMaterials;
+ NCollection_StdAllocator<OpenGl_RaytraceMaterial> > anEmptyMaterials;
Materials.swap (anEmptyMaterials);
}
-// =======================================================================
-// function : OpenGl_Node
-// purpose : Creates new empty BVH node
-// =======================================================================
-OpenGl_BVHNode::OpenGl_BVHNode()
-: myMinPoint (THE_ZERO_VEC_4F),
- myMaxPoint (THE_ZERO_VEC_4F),
- myDataRcrd (THE_ZERO_VEC_4I)
-{
- //
-}
+#ifdef HAVE_TBB
-// =======================================================================
-// function : OpenGl_Node
-// purpose : Creates new BVH node with specified data
-// =======================================================================
-OpenGl_BVHNode::OpenGl_BVHNode (const OpenGl_RTVec4f& theMinPoint,
- const OpenGl_RTVec4f& theMaxPoint,
- const OpenGl_RTVec4i& theDataRcrd)
-: myMinPoint (theMinPoint),
- myMaxPoint (theMaxPoint),
- myDataRcrd (theDataRcrd)
+struct OpenGL_BVHParallelBuilder
{
- //
-}
-
-// =======================================================================
-// function : OpenGl_Node
-// purpose : Creates new leaf BVH node with specified data
-// =======================================================================
-OpenGl_BVHNode::OpenGl_BVHNode (const OpenGl_AABB& theAABB,
- const int theBegTriangle,
- const int theEndTriangle)
-: myMinPoint (theAABB.CornerMin()),
- myMaxPoint (theAABB.CornerMax()),
- myDataRcrd (1,
- theBegTriangle,
- theEndTriangle,
- 0)
-{
- //
-}
-
-// =======================================================================
-// function : OpenGl_Node
-// purpose : Creates new leaf BVH node with specified data
-// =======================================================================
-OpenGl_BVHNode::OpenGl_BVHNode (const OpenGl_RTVec4f& theMinPoint,
- const OpenGl_RTVec4f& theMaxPoint,
- const int theBegTriangle,
- const int theEndTriangle)
-: myMinPoint (theMinPoint),
- myMaxPoint (theMaxPoint),
- myDataRcrd (1,
- theBegTriangle,
- theEndTriangle,
- 0)
-{
- //
-}
-
-// =======================================================================
-// function : CleanUp
-// purpose : Removes all tree nodes
-// =======================================================================
-void OpenGl_BVH::CleanUp()
-{
- OpenGl_RTArray4f anEmptyMinPointBuffer;
- myMinPointBuffer.swap (anEmptyMinPointBuffer);
-
- OpenGl_RTArray4f anEmptyMaxPointBuffer;
- myMaxPointBuffer.swap (anEmptyMaxPointBuffer);
-
- OpenGl_RTArray4i anEmptyDataRcrdBuffer;
- myDataRcrdBuffer.swap (anEmptyDataRcrdBuffer);
-}
-
-// =======================================================================
-// function : Node
-// purpose : Returns node with specified index
-// =======================================================================
-OpenGl_BVHNode OpenGl_BVH::Node (const int theIndex) const
-{
- return OpenGl_BVHNode (myMinPointBuffer[theIndex],
- myMaxPointBuffer[theIndex],
- myDataRcrdBuffer[theIndex]);
-}
+ BVH_ObjectSet<Standard_ShortReal, 4>* Set;
-// =======================================================================
-// function : SetNode
-// purpose : Replaces node with specified index
-// =======================================================================
-void OpenGl_BVH::SetNode (const int theIndex,
- const OpenGl_BVHNode& theNode)
-{
- if (theIndex < static_cast<int> (myMinPointBuffer.size()))
+ OpenGL_BVHParallelBuilder (BVH_ObjectSet<Standard_ShortReal, 4>* theSet)
+ : Set (theSet)
{
- myMinPointBuffer[theIndex] = theNode.myMinPoint;
- myMaxPointBuffer[theIndex] = theNode.myMaxPoint;
- myDataRcrdBuffer[theIndex] = theNode.myDataRcrd;
+ //
}
-}
-
-// =======================================================================
-// function : PushNode
-// purpose : Adds new node to the tree
-// =======================================================================
-int OpenGl_BVH::PushNode (const OpenGl_BVHNode& theNode)
-{
- myMinPointBuffer.push_back (theNode.myMinPoint);
- myMaxPointBuffer.push_back (theNode.myMaxPoint);
- myDataRcrdBuffer.push_back (theNode.myDataRcrd);
- return static_cast<int> (myDataRcrdBuffer.size() - 1);
-}
-
-// =======================================================================
-// function : OpenGl_NodeBuildTask
-// purpose : Creates new node building task
-// =======================================================================
-OpenGl_BVHNodeTask::OpenGl_BVHNodeTask()
-: NodeToBuild (0),
- BegTriangle (0),
- EndTriangle (0)
-{
- //
-}
-
-// =======================================================================
-// function : OpenGl_NodeBuildTask
-// purpose : Creates new node building task
-// =======================================================================
-OpenGl_BVHNodeTask::OpenGl_BVHNodeTask (const int theNodeToBuild,
- const int theBegTriangle,
- const int theEndTriangle)
-: NodeToBuild (theNodeToBuild),
- BegTriangle (theBegTriangle),
- EndTriangle (theEndTriangle)
-{
- //
-}
-
-// =======================================================================
-// function : OpenGl_BinnedBVHBuilder
-// purpose : Creates new binned BVH builder
-// =======================================================================
-OpenGl_BinnedBVHBuilder::OpenGl_BinnedBVHBuilder()
-: myMaxDepth (30)
-{
- //
-}
-
-// =======================================================================
-// function : ~OpenGl_BinnedBVHBuilder
-// purpose : Releases binned BVH builder
-// =======================================================================
-OpenGl_BinnedBVHBuilder::~OpenGl_BinnedBVHBuilder()
-{
- //
-}
-#define BVH_DEBUG_OUTPUT_
+ void operator() (const tbb::blocked_range<size_t>& theRange) const
+ {
+ for (size_t anObjectIdx = theRange.begin(); anObjectIdx != theRange.end(); ++anObjectIdx)
+ {
+ OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
+ Set->Objects().ChangeValue (static_cast<Standard_Integer> (anObjectIdx)).operator->());
+
+ if (aTriangleSet != NULL)
+ {
+ aTriangleSet->BVH();
+ }
+ }
+ }
+};
-#if defined( BVH_DEBUG_OUTPUT )
- #include <iostream>
#endif
// =======================================================================
-// function : ReinterpretIntAsFloat
-// purpose : Reinterprets bits of integer value as floating-point value
+// function : ProcessAcceleration
+// purpose : Performs post-processing of high-level BVH
// =======================================================================
-inline float ReinterpretIntAsFloat (int theValue)
+Standard_Boolean OpenGl_RaytraceGeometry::ProcessAcceleration()
{
- return *reinterpret_cast< float* > (&theValue);
-}
-
-// =======================================================================
-// function : Build
-// purpose : Builds BVH tree using binned SAH algorithm
-// =======================================================================
-void OpenGl_BinnedBVHBuilder::Build (OpenGl_RaytraceScene& theGeometry,
- const float theEpsilon)
-{
- CleanUp();
-
-#ifdef BVH_DEBUG_OUTPUT
- std::cout << "Start building BVH..." << std::endl;
-
- std::cout << "Triangles: " << theGeometry.Triangles.size() << std::endl;
+#ifdef BVH_PRINT_INFO
+ OSD_Timer aTimer;
#endif
+
+ MarkDirty(); // force BVH rebuilding
- if (theGeometry.Triangles.size() == 0)
- return;
-
- // Create root node with all scene triangles
- OpenGl_AABB anAABB = theGeometry.AABB;
- anAABB.CornerMin() = OpenGl_RTVec4f (anAABB.CornerMin().x() - theEpsilon,
- anAABB.CornerMin().y() - theEpsilon,
- anAABB.CornerMin().z() - theEpsilon,
- 1.0f);
- anAABB.CornerMax() = OpenGl_RTVec4f (anAABB.CornerMax().x() + theEpsilon,
- anAABB.CornerMax().y() + theEpsilon,
- anAABB.CornerMax().z() + theEpsilon,
- 1.0f);
- myTree.PushNode (OpenGl_BVHNode (anAABB, 0, static_cast<int> (theGeometry.Triangles.size() - 1)));
-
-#ifdef BVH_DEBUG_OUTPUT
- std::cout << "Push root node: [" << 0 << ", " <<
- theGeometry.Triangles.size() - 1 << "]" << std::endl;
+#ifdef BVH_PRINT_INFO
+ aTimer.Reset();
+ aTimer.Start();
#endif
- // Setup splitting task for the root node
- myNodeTasksQueue.push_back (OpenGl_BVHNodeTask (0, 0, static_cast<int> (theGeometry.Triangles.size() - 1)));
-
- // Building nodes while tasks queue is not empty
- for (int aTaskId = 0; aTaskId < static_cast<int> (myNodeTasksQueue.size()); ++aTaskId)
+#ifdef HAVE_TBB
+ // If Intel TBB is available, perform the preliminary
+ // construction of bottom-level scene BVHs
+ tbb::parallel_for (tbb::blocked_range<size_t> (0, Size()),
+ OpenGL_BVHParallelBuilder (this));
+#endif
+
+ for (Standard_Integer anObjectIdx = 0; anObjectIdx < Size(); ++anObjectIdx)
{
- BuildNode (theGeometry, aTaskId);
+ OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
+ myObjects.ChangeValue (anObjectIdx).operator->());
+
+ Standard_ASSERT_RETURN (aTriangleSet != NULL,
+ "Error! Failed to get triangulation of OpenGL element", Standard_False);
+
+ Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(),
+ "Error! Failed to update bottom-level BVH of OpenGL element", Standard_False);
}
- // Write support data to optimize traverse
- for (int aNode = 0; aNode < static_cast<int> (myTree.DataRcrdBuffer().size()); ++aNode)
- {
- OpenGl_RTVec4i aData = myTree.DataRcrdBuffer()[aNode];
- myTree.MinPointBuffer()[aNode].w() = ReinterpretIntAsFloat (aData[0] ? aData[1] : -aData[1]);
- myTree.MaxPointBuffer()[aNode].w() = ReinterpretIntAsFloat (aData[0] ? aData[2] : -aData[2]);
- }
-}
+#ifdef BVH_PRINT_INFO
+ aTimer.Stop();
-// =======================================================================
-// function : CleanUp
-// purpose : Clears previously built tree
-// =======================================================================
-void OpenGl_BinnedBVHBuilder::CleanUp()
-{
- myTree.CleanUp();
- myNodeTasksQueue.clear();
-}
+ std::cout << "Updating bottom-level BVHs (sec): " <<
+ aTimer.ElapsedTime() << std::endl;
+#endif
-// =======================================================================
-// function : SetMaxDepth
-// purpose : Sets maximum tree depth
-// =======================================================================
-void OpenGl_BinnedBVHBuilder::SetMaxDepth (const int theMaxDepth)
-{
- if (theMaxDepth > 1 && theMaxDepth < 30)
- {
- myMaxDepth = theMaxDepth - 1;
- }
-}
+#ifdef BVH_PRINT_INFO
+ aTimer.Reset();
+ aTimer.Start();
+#endif
-//! Minimum node size to split.
-static const float THE_NODE_MIN_SIZE = 1e-4f;
+ NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> > aBVH = BVH();
-// =======================================================================
-// function : BuildNode
-// purpose : Builds node using task info
-// =======================================================================
-void OpenGl_BinnedBVHBuilder::BuildNode (OpenGl_RaytraceScene& theGeometry,
- const int theTask)
-{
- OpenGl_BVHNodeTask aTask = myNodeTasksQueue[theTask];
- OpenGl_BVHNode aNode = myTree.Node (aTask.NodeToBuild);
+#ifdef BVH_PRINT_INFO
+ aTimer.Stop();
-#ifdef BVH_DEBUG_OUTPUT
- std::cout << "Build node " << aTask.NodeToBuild << ": [" <<
- aTask.BegTriangle << ", " << aTask.EndTriangle << "]" << std::endl;
+ std::cout << "Updating high-level BVH (sec): " <<
+ aTimer.ElapsedTime() << std::endl;
#endif
- OpenGl_AABB anAABB (aNode.MinPoint(), aNode.MaxPoint());
- const OpenGl_RTVec4f aNodeSize = anAABB.Size();
- const float aNodeArea = anAABB.Area();
+ Standard_ASSERT_RETURN (!aBVH.IsNull(),
+ "Error! Failed to update high-level BVH of ray-tracing scene", Standard_False);
- // Parameters for storing best split
- float aMinSplitCost = std::numeric_limits<float>::max();
+ Standard_Integer aVerticesOffset = 0;
+ Standard_Integer aElementsOffset = 0;
+ Standard_Integer aBVHNodesOffset = 0;
- int aMinSplitAxis = -1;
- int aMinSplitIndex = 0;
- int aMinSplitLftCount = 0;
- int aMinSplitRghCount = 0;
-
- OpenGl_AABB aMinSplitLftAABB;
- OpenGl_AABB aMinSplitRghAABB;
-
- // Find best split
- for (int anAxis = 0; anAxis < 3; ++anAxis)
+ for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx)
{
- if (aNodeSize[anAxis] <= THE_NODE_MIN_SIZE)
+ if (!aBVH->IsOuter (aNodeIdx))
continue;
- OpenGl_BinVector aBins (THE_NUMBER_OF_BINS);
- GetSubVolumes (theGeometry, aNode, aBins, anAxis);
+ Standard_ASSERT_RETURN (aBVH->BegPrimitive (aNodeIdx) == aBVH->EndPrimitive (aNodeIdx),
+ "Error! Invalid leaf node in high-level BVH (contains several objects)", Standard_False);
- // Choose the best split (with minimum SAH cost)
- for (int aSplit = 1; aSplit < THE_NUMBER_OF_BINS; ++aSplit)
- {
- int aLftCount = 0;
- int aRghCount = 0;
- OpenGl_AABB aLftAABB;
- OpenGl_AABB aRghAABB;
- for (int anIndex = 0; anIndex < aSplit; ++anIndex)
- {
- aLftCount += aBins[anIndex].Count;
- aLftAABB.Combine (aBins[anIndex].Volume);
- }
+ Standard_Integer anObjectIdx = aBVH->BegPrimitive (aNodeIdx);
- for (int anIndex = aSplit; anIndex < THE_NUMBER_OF_BINS; ++anIndex)
- {
- aRghCount += aBins[anIndex].Count;
- aRghAABB.Combine (aBins[anIndex].Volume);
- }
+ Standard_ASSERT_RETURN (anObjectIdx < myObjects.Size(),
+ "Error! Invalid leaf node in high-level BVH (contains out-of-range object)", Standard_False);
- // Simple SAH evaluation
- float aCost = ( aLftAABB.Area() / aNodeArea ) * aLftCount +
- ( aRghAABB.Area() / aNodeArea ) * aRghCount;
+ OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
+ myObjects.ChangeValue (anObjectIdx).operator->());
-#ifdef BVH_DEBUG_OUTPUT
- std::cout << "\t\tBin " << aSplit << ", Cost = " << aCost << std::endl;
-#endif
+ // Note: We overwrite node info record to store parameters
+ // of bottom-level BVH and triangulation of OpenGL element
- if (aCost <= aMinSplitCost)
- {
- aMinSplitCost = aCost;
- aMinSplitAxis = anAxis;
- aMinSplitIndex = aSplit;
- aMinSplitLftAABB = aLftAABB;
- aMinSplitRghAABB = aRghAABB;
- aMinSplitLftCount = aLftCount;
- aMinSplitRghCount = aRghCount;
- }
- }
- }
+ aBVH->NodeInfoBuffer().at (aNodeIdx) = BVH_Vec4i (
+ anObjectIdx + 1 /* to keep leaf flag */, aBVHNodesOffset, aVerticesOffset, aElementsOffset);
- if (aMinSplitAxis == -1)
- {
- // make outer (leaf) node
- myTree.DataRcrdBuffer()[aTask.NodeToBuild].x() = 1;
- return;
- }
-
-#ifdef BVH_DEBUG_OUTPUT
- switch (aMinSplitAxis)
- {
- case 0:
- std::cout << "\tSplit axis: X = " << aMinSplitIndex << std::endl;
- break;
- case 1:
- std::cout << "\tSplit axis: Y = " << aMinSplitIndex << std::endl;
- break;
- case 2:
- std::cout << "\tSplit axis: Z = " << aMinSplitIndex << std::endl;
- break;
+ aVerticesOffset += aTriangleSet->Vertices.size();
+ aElementsOffset += aTriangleSet->Elements.size();
+ aBVHNodesOffset += aTriangleSet->BVH()->Length();
}
-#endif
-
- int aMiddle = SplitTriangles (theGeometry, aTask.BegTriangle, aTask.EndTriangle,
- aNode, aMinSplitIndex - 1, aMinSplitAxis);
-#ifdef BVH_DEBUG_OUTPUT
- std::cout << "\tLeft child: [" << aTask.BegTriangle << ", "
- << aMiddle - 1 << "]" << std::endl;
+ return Standard_True;
+}
- std::cout << "\tRight child: [" << aMiddle << ", "
- << aTask.EndTriangle << "]" << std::endl;
-#endif
+// =======================================================================
+// function : AccelerationOffset
+// purpose : Returns offset of bottom-level BVH for given leaf node
+// =======================================================================
+Standard_Integer OpenGl_RaytraceGeometry::AccelerationOffset (Standard_Integer theNodeIdx)
+{
+ const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = BVH();
-#define BVH_SIDE_LFT 1
-#define BVH_SIDE_RGH 2
+ if (theNodeIdx >= aBVH->Length() || !aBVH->IsOuter (theNodeIdx))
+ return INVALID_OFFSET;
- // Setting up tasks for child nodes
- for (int aSide = BVH_SIDE_LFT; aSide <= BVH_SIDE_RGH; ++aSide)
- {
- OpenGl_RTVec4f aMinPoint = (aSide == BVH_SIDE_LFT)
- ? aMinSplitLftAABB.CornerMin()
- : aMinSplitRghAABB.CornerMin();
- OpenGl_RTVec4f aMaxPoint = (aSide == BVH_SIDE_LFT)
- ? aMinSplitLftAABB.CornerMax()
- : aMinSplitRghAABB.CornerMax();
-
- int aBegTriangle = (aSide == BVH_SIDE_LFT)
- ? aTask.BegTriangle
- : aMiddle;
- int aEndTriangle = (aSide == BVH_SIDE_LFT)
- ? aMiddle - 1
- : aTask.EndTriangle;
-
- OpenGl_BVHNode aChild (aMinPoint, aMaxPoint, aBegTriangle, aEndTriangle);
- aChild.SetLevel (aNode.Level() + 1);
-
- // Check to see if child node must be split
- const int aNbTriangles = (aSide == BVH_SIDE_LFT)
- ? aMinSplitLftCount
- : aMinSplitRghCount;
-
- const int isChildALeaf = (aNbTriangles <= THE_MAX_LEAF_TRIANGLES) || (aNode.Level() >= myMaxDepth);
- if (isChildALeaf)
- aChild.SetOuter();
- else
- aChild.SetInner();
-
- const int aChildIndex = myTree.PushNode (aChild);
-
- // Modify parent node
- myTree.DataRcrdBuffer()[aTask.NodeToBuild].x() = 0; // inner node flag
- if (aSide == BVH_SIDE_LFT)
- myTree.DataRcrdBuffer()[aTask.NodeToBuild].y() = aChildIndex; // left child
- else
- myTree.DataRcrdBuffer()[aTask.NodeToBuild].z() = aChildIndex; // right child
-
- // Make new building task
- if (!isChildALeaf)
- myNodeTasksQueue.push_back (OpenGl_BVHNodeTask (aChildIndex, aBegTriangle, aEndTriangle));
- }
+ return aBVH->NodeInfoBuffer().at (theNodeIdx).y();
}
// =======================================================================
-// function : SplitTriangles
-// purpose : Splits node triangles into two intervals for child nodes
+// function : VerticesOffset
+// purpose : Returns offset of triangulation vertices for given leaf node
// =======================================================================
-int OpenGl_BinnedBVHBuilder::SplitTriangles (OpenGl_RaytraceScene& theGeometry,
- const int theBegTriangle,
- const int theEndTriangle,
- OpenGl_BVHNode& theNode,
- int theBin,
- const int theAxis)
+Standard_Integer OpenGl_RaytraceGeometry::VerticesOffset (Standard_Integer theNodeIdx)
{
- int aLftIndex (theBegTriangle);
- int aRghIndex (theEndTriangle);
-
- const float aMin = theNode.MinPoint()[theAxis];
- const float aMax = theNode.MaxPoint()[theAxis];
+ const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = BVH();
- const float aStep = (aMax - aMin) / THE_NUMBER_OF_BINS;
+ if (theNodeIdx >= aBVH->Length() || !aBVH->IsOuter (theNodeIdx))
+ return INVALID_OFFSET;
- do
- {
- while ((int )floorf ((theGeometry.CenterAxis (aLftIndex, theAxis) - aMin) / aStep) <= theBin
- && aLftIndex < theEndTriangle)
- {
- ++aLftIndex;
- }
- while ((int )floorf ((theGeometry.CenterAxis (aRghIndex, theAxis) - aMin) / aStep) > theBin
- && aRghIndex > theBegTriangle)
- {
- --aRghIndex;
- }
+ return aBVH->NodeInfoBuffer().at (theNodeIdx).z();
+}
- if (aLftIndex <= aRghIndex)
- {
- if (aLftIndex != aRghIndex)
- {
- OpenGl_RTVec4i aLftTrg = theGeometry.Triangles[aLftIndex];
- OpenGl_RTVec4i aRghTrg = theGeometry.Triangles[aRghIndex];
- theGeometry.Triangles[aLftIndex] = aRghTrg;
- theGeometry.Triangles[aRghIndex] = aLftTrg;
- }
+// =======================================================================
+// function : ElementsOffset
+// purpose : Returns offset of triangulation elements for given leaf node
+// =======================================================================
+Standard_Integer OpenGl_RaytraceGeometry::ElementsOffset (Standard_Integer theNodeIdx)
+{
+ const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = BVH();
- aLftIndex++; aRghIndex--;
- }
- } while (aLftIndex <= aRghIndex);
+ if (theNodeIdx >= aBVH->Length() || !aBVH->IsOuter (theNodeIdx))
+ return INVALID_OFFSET;
- return aLftIndex;
+ return aBVH->NodeInfoBuffer().at (theNodeIdx).w();
}
// =======================================================================
-// function : GetSubVolumes
-// purpose : Arranges node triangles into bins
+// function : TriangleSet
+// purpose : Returns triangulation data for given leaf node
// =======================================================================
-void OpenGl_BinnedBVHBuilder::GetSubVolumes (OpenGl_RaytraceScene& theGeometry,
- const OpenGl_BVHNode& theNode,
- OpenGl_BinVector& theBins,
- const int theAxis)
+OpenGl_TriangleSet* OpenGl_RaytraceGeometry::TriangleSet (Standard_Integer theNodeIdx)
{
- const float aMin = theNode.MinPoint()[theAxis];
- const float aMax = theNode.MaxPoint()[theAxis];
+ const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = BVH();
- const float aStep = (aMax - aMin) / THE_NUMBER_OF_BINS;
+ if (theNodeIdx >= aBVH->Length() || !aBVH->IsOuter (theNodeIdx))
+ return NULL;
- for (int aTri = theNode.BegTriangle(); aTri <= theNode.EndTriangle(); ++aTri)
- {
- float aCenter = theGeometry.CenterAxis (aTri, theAxis);
- int aBinIndex = (int )floorf ((aCenter - aMin) * ( 1.0f / aStep));
- if (aBinIndex < 0)
- {
- aBinIndex = 0;
- }
- else if (aBinIndex >= THE_NUMBER_OF_BINS)
- {
- aBinIndex = THE_NUMBER_OF_BINS - 1;
- }
-
- theBins[aBinIndex].Count++;
- theBins[aBinIndex].Volume.Combine (theGeometry.Box (aTri));
- }
+ if (aBVH->NodeInfoBuffer().at (theNodeIdx).x() > myObjects.Size())
+ return NULL;
+
+ return dynamic_cast<OpenGl_TriangleSet*> (myObjects.ChangeValue (
+ aBVH->NodeInfoBuffer().at (theNodeIdx).x() - 1).operator->());
}
namespace OpenGl_Raytrace
#ifdef HAVE_OPENCL
-#include <OpenGl_AABB.hxx>
-#include <OpenGl_Structure.hxx>
+#include <BVH_Geometry.hxx>
+#include <BVH_Triangulation.hxx>
+#include <NCollection_StdAllocator.hxx>
#include <OpenGl_PrimitiveArray.hxx>
+#include <OpenGl_Structure.hxx>
namespace OpenGl_Raytrace
{
public:
//! Ambient reflection coefficient.
- OpenGl_RTVec4f Ambient;
+ BVH_Vec4f Ambient;
//! Diffuse reflection coefficient.
- OpenGl_RTVec4f Diffuse;
+ BVH_Vec4f Diffuse;
//! Glossy reflection coefficient.
- OpenGl_RTVec4f Specular;
+ BVH_Vec4f Specular;
//! Material emission.
- OpenGl_RTVec4f Emission;
+ BVH_Vec4f Emission;
//! Specular reflection coefficient.
- OpenGl_RTVec4f Reflection;
+ BVH_Vec4f Reflection;
//! Specular refraction coefficient.
- OpenGl_RTVec4f Refraction;
+ BVH_Vec4f Refraction;
//! Material transparency.
- OpenGl_RTVec4f Transparency;
+ BVH_Vec4f Transparency;
public:
OpenGl_RaytraceMaterial();
//! Creates new material with specified properties.
- OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient,
- const OpenGl_RTVec4f& theDiffuse,
- const OpenGl_RTVec4f& theSpecular);
+ OpenGl_RaytraceMaterial (const BVH_Vec4f& theAmbient,
+ const BVH_Vec4f& theDiffuse,
+ const BVH_Vec4f& theSpecular);
//! Creates new material with specified properties.
- OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient,
- const OpenGl_RTVec4f& theDiffuse,
- const OpenGl_RTVec4f& theSpecular,
- const OpenGl_RTVec4f& theEmission,
- const OpenGl_RTVec4f& theTranspar);
+ OpenGl_RaytraceMaterial (const BVH_Vec4f& theAmbient,
+ const BVH_Vec4f& theDiffuse,
+ const BVH_Vec4f& theSpecular,
+ const BVH_Vec4f& theEmission,
+ const BVH_Vec4f& theTranspar);
//! Creates new material with specified properties.
- OpenGl_RaytraceMaterial (const OpenGl_RTVec4f& theAmbient,
- const OpenGl_RTVec4f& theDiffuse,
- const OpenGl_RTVec4f& theSpecular,
- const OpenGl_RTVec4f& theEmission,
- const OpenGl_RTVec4f& theTranspar,
- const OpenGl_RTVec4f& theReflection,
- const OpenGl_RTVec4f& theRefraction);
+ OpenGl_RaytraceMaterial (const BVH_Vec4f& theAmbient,
+ const BVH_Vec4f& theDiffuse,
+ const BVH_Vec4f& theSpecular,
+ const BVH_Vec4f& theEmission,
+ const BVH_Vec4f& theTranspar,
+ const BVH_Vec4f& theReflection,
+ const BVH_Vec4f& theRefraction);
//! Returns packed (serialized) representation of material.
- const float* Packed() { return reinterpret_cast<float*> (this); }
+ const Standard_ShortReal* Packed()
+ {
+ return reinterpret_cast<Standard_ShortReal*> (this);
+ }
};
//! Stores properties of OpenGL light source.
{
public:
- //! 'Ambient' intensity.
- OpenGl_RTVec4f Ambient;
-
- //! 'Diffuse' intensity.
- OpenGl_RTVec4f Diffuse;
+ //! Diffuse intensity (in terms of OpenGL).
+ BVH_Vec4f Diffuse;
//! Position of light source (in terms of OpenGL).
- OpenGl_RTVec4f Position;
-
+ BVH_Vec4f Position;
public:
//! Creates new light source.
- OpenGl_RaytraceLight (const OpenGl_RTVec4f& theAmbient);
-
- //! Creates new light source.
- OpenGl_RaytraceLight (const OpenGl_RTVec4f& theDiffuse,
- const OpenGl_RTVec4f& thePosition);
+ OpenGl_RaytraceLight (const BVH_Vec4f& theDiffuse,
+ const BVH_Vec4f& thePosition);
//! Returns packed (serialized) representation of light source.
- const float* Packed() { return reinterpret_cast<float*> (this); }
+ const Standard_ShortReal* Packed()
+ {
+ return reinterpret_cast<Standard_ShortReal*> (this);
+ }
};
-//! Stores scene geometry data.
-struct OpenGl_RaytraceScene
+//! Triangulation of single OpenGL primitive array.
+class OpenGl_TriangleSet : public BVH_Triangulation<Standard_ShortReal, 4>
{
- //! AABB of 3D scene.
- OpenGl_AABB AABB;
+public:
//! Array of vertex normals.
- OpenGl_RTArray4f Normals;
-
- //! Array of vertex coordinates.
- OpenGl_RTArray4f Vertices;
-
- //! Array of scene triangles.
- OpenGl_RTArray4i Triangles;
-
- //! Array of 'front' material properties.
- std::vector<OpenGl_RaytraceMaterial,
- NCollection_StdAllocator<OpenGl_RaytraceMaterial> > Materials;
-
- //! Array of properties of light sources.
- std::vector<OpenGl_RaytraceLight,
- NCollection_StdAllocator<OpenGl_RaytraceLight> > LightSources;
-
- //! Clears all scene geometry and material data.
- void Clear();
-
- //! Returns AABB of specified triangle.
- OpenGl_AABB Box (const int theTriangle) const;
-
- //! Returns centroid of specified triangle.
- OpenGl_RTVec4f Center (const int theTriangle) const;
-
- //! Returns centroid coordinate for specified axis.
- float CenterAxis (const int theTriangle, const int theAxis) const;
-};
-
-//! Stores parameters of BVH tree node.
-class OpenGl_BVHNode
-{
- friend class OpenGl_BVH;
+ BVH_Array4f Normals;
public:
- //! Creates new empty BVH node.
- OpenGl_BVHNode();
-
- //! Creates new BVH node with specified data.
- OpenGl_BVHNode (const OpenGl_RTVec4f& theMinPoint,
- const OpenGl_RTVec4f& theMaxPoint,
- const OpenGl_RTVec4i& theDataRcrd);
-
- //! Creates new leaf BVH node with specified data.
- OpenGl_BVHNode (const OpenGl_RTVec4f& theMinPoint,
- const OpenGl_RTVec4f& theMaxPoint,
- const int theBegTriangle,
- const int theEndTriangle);
-
- //! Creates new leaf BVH node with specified data.
- OpenGl_BVHNode (const OpenGl_AABB& theAABB,
- const int theBegTriangle,
- const int theEndTriangle);
-
- //! Returns minimum point of node's AABB.
- OpenGl_RTVec4f& MinPoint() { return myMinPoint; }
- //! Returns maximum point of node's AABB.
- OpenGl_RTVec4f& MaxPoint() { return myMaxPoint; }
-
- //! Returns minimum point of node's AABB.
- const OpenGl_RTVec4f& MinPoint() const { return myMinPoint; }
- //! Returns maximum point of node's AABB.
- const OpenGl_RTVec4f& MaxPoint() const { return myMaxPoint; }
-
- //! Returns index of left child of inner node.
- int LeftChild() const { return myDataRcrd.y(); }
- //! Sets index of left child of inner node.
- void SetLeftChild (int theChild) { myDataRcrd.y() = theChild; }
-
- //! Returns index of right child of inner node.
- int RightChild() const { return myDataRcrd.z(); }
- //! Sets index of right child of inner node.
- void SetRightChild (int theChild) { myDataRcrd.z() = theChild; }
-
- //! Returns index of begin triangle of leaf node.
- int BegTriangle() const { return myDataRcrd.y(); }
- //! Sets index of begin triangle of leaf node.
- void SetBegTriangle (int theIndex) { myDataRcrd.y() = theIndex; }
-
- //! Returns index of end triangle of leaf node.
- int EndTriangle() const { return myDataRcrd.z(); }
- //! Sets index of end triangle of leaf node.
- void SetEndTriangle (int theIndex) { myDataRcrd.z() = theIndex; }
-
- //! Returns level of the node in BVH tree.
- int Level() const { return myDataRcrd.w(); }
- //! Sets level of the node in BVH tree.
- void SetLevel (int theLevel) { myDataRcrd.w() = theLevel; }
-
- //! Is node a leaf (outer)?
- bool IsOuter() const { return myDataRcrd.x() == 1; }
-
- //! Sets node type to 'outer'.
- void SetOuter() { myDataRcrd.x() = 1; }
- //! Sets node type to 'inner'.
- void SetInner() { myDataRcrd.x() = 0; }
-
-private:
-
- //! Minimum point of node's bounding box.
- OpenGl_RTVec4f myMinPoint;
- //! Maximum point of node's bounding box.
- OpenGl_RTVec4f myMaxPoint;
-
- //! Data vector (stores data fields of the node).
- OpenGl_RTVec4i myDataRcrd;
+ //! Creates new OpenGL element triangulation.
+ OpenGl_TriangleSet()
+ {
+ //
+ }
+
+ //! Releases resources of OpenGL element triangulation.
+ ~OpenGl_TriangleSet()
+ {
+ //
+ }
};
-//! Stores parameters of BVH tree.
-class OpenGl_BVH
+//! Stores geometry of ray-tracing scene.
+class OpenGl_RaytraceGeometry : public BVH_Geometry<Standard_ShortReal, 4>
{
public:
- //! Removes all tree nodes.
- void CleanUp();
-
- //! Adds new node to the tree.
- int PushNode (const OpenGl_BVHNode& theNode);
+ //! Value of invalid offset to return in case of errors.
+ static const Standard_Integer INVALID_OFFSET = -1;
- //! Returns node with specified index.
- OpenGl_BVHNode Node (const int theIndex) const;
-
- //! Replaces node with specified index by the new one.
- void SetNode (const int theIndex, const OpenGl_BVHNode& theNode);
-
- //! Returns array of node min points.
- OpenGl_RTArray4f& MinPointBuffer() { return myMinPointBuffer; }
- //! Returns array of node max points.
- OpenGl_RTArray4f& MaxPointBuffer() { return myMaxPointBuffer; }
- //! Returns array of node data records.
- OpenGl_RTArray4i& DataRcrdBuffer() { return myDataRcrdBuffer; }
-
-private:
-
- //! Array of min points of BVH nodes.
- OpenGl_RTArray4f myMinPointBuffer;
- //! Array of max points of BVH nodes.
- OpenGl_RTArray4f myMaxPointBuffer;
- //! Array of data vectors of BVH nodes.
- OpenGl_RTArray4i myDataRcrdBuffer;
-};
-
-//! Stores parameters of single node bin (slice of AABB).
-struct OpenGl_BVHBin
-{
- //! Creates new node bin.
- OpenGl_BVHBin(): Count (0) { }
-
- //! Number of primitives in the bin.
- int Count;
-
- //! AABB of the bin.
- OpenGl_AABB Volume;
-};
-
-//! Node building task.
-struct OpenGl_BVHNodeTask
-{
- //! Creates new node building task.
- OpenGl_BVHNodeTask();
-
- //! Creates new node building task.
- OpenGl_BVHNodeTask (const int theNodeToBuild,
- const int theBegTriangle,
- const int theEndTriangle);
-
- //! Index of building tree node.
- int NodeToBuild;
- //! Index of start node triangle.
- int BegTriangle;
- //! Index of final node triangle.
- int EndTriangle;
-};
-
-//! The array of bins of BVH tree node.
-typedef std::vector<OpenGl_BVHBin,
- NCollection_StdAllocator<OpenGl_BVHBin> > OpenGl_BinVector;
-
-//! Binned SAH-based BVH builder.
-class OpenGl_BinnedBVHBuilder
-{
public:
- //! Creates new binned BVH builder.
- OpenGl_BinnedBVHBuilder();
-
- //! Releases binned BVH builder.
- ~OpenGl_BinnedBVHBuilder();
+ //! Array of properties of light sources.
+ std::vector<OpenGl_RaytraceLight,
+ NCollection_StdAllocator<OpenGl_RaytraceLight> > Sources;
- //! Builds BVH tree using binned SAH algorithm.
- void Build (OpenGl_RaytraceScene& theGeometry, const float theEpsilon = 1e-3f);
+ //! Array of 'front' material properties.
+ std::vector<OpenGl_RaytraceMaterial,
+ NCollection_StdAllocator<OpenGl_RaytraceMaterial> > Materials;
- //! Sets maximum tree depth.
- void SetMaxDepth (const int theMaxDepth);
+ //! Global ambient from all light sources.
+ BVH_Vec4f GlobalAmbient;
- //! Clears previously constructed BVH tree.
- void CleanUp();
+public:
- //! Return constructed BVH tree.
- OpenGl_BVH& Tree() { return myTree; }
+ //! Creates uninitialized ray-tracing geometry.
+ OpenGl_RaytraceGeometry()
+ {
+ //
+ }
-private:
+ //! Releases resources of ray-tracing geometry.
+ ~OpenGl_RaytraceGeometry()
+ {
+ //
+ }
- //! Builds node using task info.
- void BuildNode (OpenGl_RaytraceScene& theGeometry, const int theTask);
+ //! Clears ray-tracing geometry.
+ void Clear();
- //! Arranges node triangles into bins.
- void GetSubVolumes (OpenGl_RaytraceScene& theGeometry, const OpenGl_BVHNode& theNode,
- OpenGl_BinVector& theBins, const int theAxis);
+public:
- //! Splits node triangles into two intervals for child nodes.
- int SplitTriangles (OpenGl_RaytraceScene& theGeometry, const int theFirst, const int theLast,
- OpenGl_BVHNode& theNode, int theBin, const int theAxis);
+ //! Performs post-processing of high-level scene BVH.
+ Standard_Boolean ProcessAcceleration();
-private:
+ //! Returns offset of bottom-level BVH for given leaf node.
+ //! If the node index is not valid the function returns -1.
+ //! @note Can be used after processing acceleration structure.
+ Standard_Integer AccelerationOffset (Standard_Integer theNodeIdx);
- //! Queue of node building tasks.
- std::vector<OpenGl_BVHNodeTask> myNodeTasksQueue;
+ //! Returns offset of triangulation vertices for given leaf node.
+ //! If the node index is not valid the function returns -1.
+ //! @note Can be used after processing acceleration structure.
+ Standard_Integer VerticesOffset (Standard_Integer theNodeIdx);
- //! Builded BVH tree.
- OpenGl_BVH myTree;
+ //! Returns offset of triangulation elements for given leaf node.
+ //! If the node index is not valid the function returns -1.
+ //! @note Can be used after processing acceleration structure.
+ Standard_Integer ElementsOffset (Standard_Integer theNodeIdx);
- //! Maximum depth of BVH tree.
- int myMaxDepth;
+ //! Returns triangulation data for given leaf node.
+ //! 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);
};
#endif
//! Updates environment map for ray-tracing.
Standard_Boolean UpdateRaytraceEnvironmentMap();
-
+
//! Adds OpenGL structure to ray-traced scene geometry.
- Standard_Boolean AddRaytraceStructure (const OpenGl_Structure* theStruct,
- const float* theTrans, std::set<const OpenGl_Structure*>& theElements);
+ Standard_Boolean AddRaytraceStructure (const OpenGl_Structure* theStructure,
+ const Standard_ShortReal* theTransform, std::set<const OpenGl_Structure*>& theElements);
//! Adds OpenGL primitive array to ray-traced scene geometry.
- Standard_Boolean AddRaytracePrimitiveArray (
- const CALL_DEF_PARRAY* theArray, int theMatID, const float* theTrans);
+ OpenGl_TriangleSet* AddRaytracePrimitiveArray (
+ const CALL_DEF_PARRAY* theArray, int theMatID, const Standard_ShortReal* theTrans);
//! Adds vertex indices from OpenGL primitive array to ray-traced scene geometry.
- Standard_Boolean AddRaytraceVertexIndices (const CALL_DEF_PARRAY* theArray,
- int theFirstVert, int theVertOffset, int theVertNum, int theMatID);
+ Standard_Boolean AddRaytraceVertexIndices (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID);
//! Adds OpenGL triangle array to ray-traced scene geometry.
- Standard_Boolean AddRaytraceTriangleArray (const CALL_DEF_PARRAY* theArray,
- int theFirstVert, int theVertOffset, int theVertNum, int theMatID);
+ Standard_Boolean AddRaytraceTriangleArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID);
//! Adds OpenGL triangle fan array to ray-traced scene geometry.
- Standard_Boolean AddRaytraceTriangleFanArray (const CALL_DEF_PARRAY* theArray,
- int theFirstVert, int theVertOffset, int theVertNum, int theMatID);
+ Standard_Boolean AddRaytraceTriangleFanArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID);
- //! Adds OpenGL triangle fan array to ray-traced scene geometry.
- Standard_Boolean AddRaytraceTriangleStripArray (const CALL_DEF_PARRAY* theArray,
- int theFirstVert, int theVertOffset, int theVertNum, int theMatID);
+ //! Adds OpenGL triangle strip array to ray-traced scene geometry.
+ Standard_Boolean AddRaytraceTriangleStripArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID);
//! Adds OpenGL quadrangle array to ray-traced scene geometry.
- Standard_Boolean AddRaytraceQuadrangleArray (const CALL_DEF_PARRAY* theArray,
- int theFirstVert, int theVertOffset, int theVertNum, int theMatID);
+ Standard_Boolean AddRaytraceQuadrangleArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID);
//! Adds OpenGL quadrangle strip array to ray-traced scene geometry.
- Standard_Boolean AddRaytraceQuadrangleStripArray (const CALL_DEF_PARRAY* theArray,
- int theFirstVert, int theVertOffset, int theVertNum, int theMatID);
+ Standard_Boolean AddRaytraceQuadrangleStripArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID);
//! Adds OpenGL polygon array to ray-traced scene geometry.
- Standard_Boolean AddRaytracePolygonArray (const CALL_DEF_PARRAY* theArray,
- int theFirstVert, int theVertOffset, int theVertNum, int theMatID);
+ 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();
Standard_Boolean WriteRaytraceSceneToDevice();
//! Runs OpenCL ray-tracing kernels.
- Standard_Boolean RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView,
+ Standard_Boolean RunRaytraceOpenCLKernelsOld (const Graphic3d_CView& theCView,
const GLfloat theOrigins[16],
const GLfloat theDirects[16],
const int theSizeX,
const int theSizeY);
+ //! 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);
+
//! Redraws the window using OpenCL ray tracing.
Standard_Boolean Raytrace (const Graphic3d_CView& theCView,
const int theSizeX, const int theSizeY, const Tint theToSwap);
Standard_Boolean myIsRaytraceDataValid;
//! Is geometry data musty be updated?
Standard_Boolean myToUpdateRaytraceData;
+
//! 3D scene geometry data for ray-tracing.
- OpenGl_RaytraceScene myRaytraceSceneData;
+ OpenGl_RaytraceGeometry myRaytraceGeometry;
//! Radius of bounding sphere of the scene.
- float myRaytraceSceneRadius;
+ Standard_ShortReal myRaytraceSceneRadius;
//! Scene epsilon to prevent self-intersections.
- float myRaytraceSceneEpsilon;
-
- //! Binned SAH-based BVH builder.
- OpenGl_BinnedBVHBuilder myBVHBuilder;
+ Standard_ShortReal myRaytraceSceneEpsilon;
//! OpenCL context.
cl_context myComputeContext;
//! OpenCL command queue.
- cl_command_queue myRaytraceQueue;
+ cl_command_queue myComputeQueue;
//! OpenCL computing program.
cl_program myRaytraceProgram;
//! OpenCL ray-tracing render kernel.
//! OpenGL texture to store anti-aliasing result.
Handle(OpenGl_Texture) myRaytraceOutputTextureAA;
- //! OpenCL buffer of vertex normals.
- cl_mem myRaytraceNormalBuffer;
- //! OpenCL buffer of vertex coordinates.
- cl_mem myRaytraceVertexBuffer;
- //! OpenCL buffer of indices of triangle vertices.
- cl_mem myRaytraceTriangleBuffer;
-
- //! OpenCL buffer of minimum points of BVH nodes.
- cl_mem myRaytraceNodeMinPointBuffer;
- //! OpenCL buffer of maximum points of BVH nodes.
- cl_mem myRaytraceNodeMaxPointBuffer;
- //! OpenCL buffer of data records of BVH nodes.
- cl_mem myRaytraceNodeDataRcrdBuffer;
-
//! 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;
+
//! State of OpenGL view.
Standard_Size myViewModificationStatus;
-
//! State of OpenGL layer list.
Standard_Size myLayersModificationStatus;
- //! State of OpenGL elements reflected to ray-tracing.
+ //! State of OpenGL structures reflected to ray-tracing.
std::map<const OpenGl_Structure*, Standard_Size> myStructureStates;
#endif // HAVE_OPENCL
#include <OpenGl_Texture.hxx>
#include <OpenGl_View.hxx>
#include <OpenGl_Workspace.hxx>
+#include <Standard_Assert.hxx>
using namespace OpenGl_Raytrace;
// function : MatVecMult
// purpose : Multiples 4x4 matrix by 4D vector
// =======================================================================
-template< typename T >
-OpenGl_RTVec4f MatVecMult (const T m[16], const OpenGl_RTVec4f& v)
+template<typename T>
+BVH_Vec4f MatVecMult (const T m[16], const BVH_Vec4f& v)
{
- return OpenGl_RTVec4f (
+ return BVH_Vec4f (
static_cast<float> (m[ 0] * v.x() + m[ 4] * v.y() +
m[ 8] * v.z() + m[12] * v.w()),
static_cast<float> (m[ 1] * v.x() + m[ 5] * v.y() +
aSizeY = (myView->TextureEnv()->SizeY() <= 0) ? 1 : myView->TextureEnv()->SizeY();
}
- cl_image_format aImageFormat;
+ cl_image_format anImageFormat;
- aImageFormat.image_channel_order = CL_RGBA;
- aImageFormat.image_channel_data_type = CL_FLOAT;
+ anImageFormat.image_channel_order = CL_RGBA;
+ anImageFormat.image_channel_data_type = CL_FLOAT;
- myRaytraceEnvironment = clCreateImage2D (myComputeContext, CL_MEM_READ_ONLY,
- &aImageFormat, aSizeX, aSizeY, 0,
- NULL, &anError);
+ myRaytraceEnvironment = clCreateImage2D (myComputeContext,
+ CL_MEM_READ_ONLY, &anImageFormat, aSizeX, aSizeY, 0, NULL, &anError);
cl_float* aPixelData = new cl_float[aSizeX * aSizeY * 4];
aSizeY,
1 };
- anError |= clEnqueueWriteImage (myRaytraceQueue, myRaytraceEnvironment, CL_TRUE,
- anImageOffset, anImageRegion, 0, 0, aPixelData,
- 0, NULL, NULL);
+ 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;
if (!theCheck)
{
- myRaytraceSceneData.Clear();
+ myRaytraceGeometry.Clear();
myIsRaytraceDataValid = Standard_False;
}
}
}
- float* aTransform (NULL);
+ Standard_ShortReal* aTransform (NULL);
// The set of processed structures (reflected to ray-tracing)
// This set is used to remove out-of-date records from the
if (aStructure->Transformation()->mat != NULL)
{
if (aTransform == NULL)
- aTransform = new float[16];
+ aTransform = new Standard_ShortReal[16];
for (Standard_Integer i = 0; i < 4; ++i)
for (Standard_Integer j = 0; j < 4; ++j)
// Actualize OpenGL layer list state
myLayersModificationStatus = myView->LayerList().ModificationState();
+ // Rebuild bottom-level and high-level BVHs
+ myRaytraceGeometry.ProcessAcceleration();
-#ifdef RAY_TRACE_PRINT_INFO
- OSD_Timer aTimer;
- aTimer.Start();
-#endif
-
- myBVHBuilder.Build (myRaytraceSceneData);
-
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << " Build time: " << aTimer.ElapsedTime() << " for "
- << myRaytraceSceneData.Triangles.size() / 1000 << "K triangles" << std::endl;
-#endif
-
- const float aScaleFactor = 1.5f;
+ const Standard_ShortReal aMinRadius = Max (fabs (myRaytraceGeometry.Box().CornerMin().x()), Max (
+ fabs (myRaytraceGeometry.Box().CornerMin().y()), fabs (myRaytraceGeometry.Box().CornerMin().z())));
+ const Standard_ShortReal aMaxRadius = Max (fabs (myRaytraceGeometry.Box().CornerMax().x()), Max (
+ fabs (myRaytraceGeometry.Box().CornerMax().y()), fabs (myRaytraceGeometry.Box().CornerMax().z())));
- myRaytraceSceneRadius = aScaleFactor *
- Max ( Max (fabsf (myRaytraceSceneData.AABB.CornerMin().x()),
- Max (fabsf (myRaytraceSceneData.AABB.CornerMin().y()),
- fabsf (myRaytraceSceneData.AABB.CornerMin().z()))),
- Max (fabsf (myRaytraceSceneData.AABB.CornerMax().x()),
- Max (fabsf (myRaytraceSceneData.AABB.CornerMax().y()),
- fabsf (myRaytraceSceneData.AABB.CornerMax().z()))) );
+ myRaytraceSceneRadius = 2.f /* scale factor */ * Max (aMinRadius, aMaxRadius);
- myRaytraceSceneEpsilon = Max (1e-4f, myRaytraceSceneRadius * 1e-4f);
+ myRaytraceSceneEpsilon = Max (1e-4f,
+ myRaytraceGeometry.Box().Size().Length() * 1e-4f);
return WriteRaytraceSceneToDevice();
}
// =======================================================================
// function : CheckRaytraceStructure
-// purpose : Adds OpenGL structure to ray-traced scene geometry
+// purpose : Checks to see if the structure is modified
// =======================================================================
Standard_Boolean OpenGl_Workspace::CheckRaytraceStructure (const OpenGl_Structure* theStructure)
{
// function : CreateMaterial
// purpose : Creates ray-tracing material properties
// =======================================================================
-void CreateMaterial (const OPENGL_SURF_PROP& theProp,
- OpenGl_RaytraceMaterial& theMaterial)
+void CreateMaterial (const OPENGL_SURF_PROP& theProp, OpenGl_RaytraceMaterial& theMaterial)
{
const float* aSrcAmb = theProp.isphysic ? theProp.ambcol.rgb : theProp.matcol.rgb;
- theMaterial.Ambient = OpenGl_RTVec4f (aSrcAmb[0] * theProp.amb,
- aSrcAmb[1] * theProp.amb,
- aSrcAmb[2] * theProp.amb,
- 1.0f);
+ theMaterial.Ambient = BVH_Vec4f (aSrcAmb[0] * theProp.amb,
+ aSrcAmb[1] * theProp.amb,
+ aSrcAmb[2] * theProp.amb,
+ 1.0f);
const float* aSrcDif = theProp.isphysic ? theProp.difcol.rgb : theProp.matcol.rgb;
- theMaterial.Diffuse = OpenGl_RTVec4f (aSrcDif[0] * theProp.diff,
- aSrcDif[1] * theProp.diff,
- aSrcDif[2] * theProp.diff,
- 1.0f);
+ theMaterial.Diffuse = BVH_Vec4f (aSrcDif[0] * theProp.diff,
+ aSrcDif[1] * theProp.diff,
+ aSrcDif[2] * theProp.diff,
+ 1.0f);
const float aDefSpecCol[4] = {1.0f, 1.0f, 1.0f, 1.0f};
const float* aSrcSpe = theProp.isphysic ? theProp.speccol.rgb : aDefSpecCol;
- theMaterial.Specular = OpenGl_RTVec4f (aSrcSpe[0] * theProp.spec,
- aSrcSpe[1] * theProp.spec,
- aSrcSpe[2] * theProp.spec,
- theProp.shine);
+ theMaterial.Specular = BVH_Vec4f (aSrcSpe[0] * theProp.spec,
+ aSrcSpe[1] * theProp.spec,
+ aSrcSpe[2] * theProp.spec,
+ theProp.shine);
const float* aSrcEms = theProp.isphysic ? theProp.emscol.rgb : theProp.matcol.rgb;
- theMaterial.Emission = OpenGl_RTVec4f (aSrcEms[0] * theProp.emsv,
- aSrcEms[1] * theProp.emsv,
- aSrcEms[2] * theProp.emsv,
- 1.0f);
+ theMaterial.Emission = BVH_Vec4f (aSrcEms[0] * theProp.emsv,
+ aSrcEms[1] * theProp.emsv,
+ aSrcEms[2] * theProp.emsv,
+ 1.0f);
// Note: Here we use sub-linear transparency function
// to produce realistic-looking transparency effect
- theMaterial.Transparency = OpenGl_RTVec4f (powf (theProp.trans, 0.75f),
- 1.f - theProp.trans,
- 1.f,
- 1.f);
+ theMaterial.Transparency = BVH_Vec4f (powf (theProp.trans, 0.75f),
+ 1.f - theProp.trans,
+ 1.f,
+ 1.f);
const float aMaxRefl = Max (theMaterial.Diffuse.x() + theMaterial.Specular.x(),
Max (theMaterial.Diffuse.y() + theMaterial.Specular.y(),
const float aReflectionScale = 0.75f / aMaxRefl;
- theMaterial.Reflection = OpenGl_RTVec4f (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,
+ theProp.speccol.rgb[1] * theProp.spec,
+ theProp.speccol.rgb[2] * theProp.spec,
+ 0.f) * aReflectionScale;
}
// =======================================================================
// function : AddRaytraceStructure
// purpose : Adds OpenGL structure to ray-traced scene geometry
// =======================================================================
-Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure* theStructure,
- const float* theTransform,
- std::set<const OpenGl_Structure*>& theElements)
+Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure* theStructure,
+ const Standard_ShortReal* theTransform, std::set<const OpenGl_Structure*>& theElements)
{
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Add Structure" << std::endl;
-#endif
-
theElements.insert (theStructure);
if (!theStructure->IsVisible())
if (theStructure->AspectFace() != NULL)
{
- aStructMatID = static_cast<Standard_Integer> (myRaytraceSceneData.Materials.size());
+ aStructMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
OpenGl_RaytraceMaterial aStructMaterial;
CreateMaterial (theStructure->AspectFace()->IntFront(), aStructMaterial);
- myRaytraceSceneData.Materials.push_back (aStructMaterial);
+ myRaytraceGeometry.Materials.push_back (aStructMaterial);
}
- OpenGl_ListOfGroup::Iterator anItg (theStructure->Groups());
-
- while (anItg.More())
+ for (OpenGl_ListOfGroup::Iterator anItg (theStructure->Groups()); anItg.More(); anItg.Next())
{
// Get group material
Standard_Integer aGroupMatID = -1;
if (anItg.Value()->AspectFace() != NULL)
{
- aGroupMatID = static_cast<Standard_Integer> (myRaytraceSceneData.Materials.size());
+ aGroupMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
OpenGl_RaytraceMaterial aGroupMaterial;
CreateMaterial (anItg.Value()->AspectFace()->IntFront(), aGroupMaterial);
- myRaytraceSceneData.Materials.push_back (aGroupMaterial);
+ myRaytraceGeometry.Materials.push_back (aGroupMaterial);
}
Standard_Integer aMatID = aGroupMatID < 0 ? aStructMatID : aGroupMatID;
- if (aStructMatID < 0 && aGroupMatID < 0)
+ if (aMatID < 0)
{
- aMatID = static_cast<Standard_Integer> (myRaytraceSceneData.Materials.size());
+ aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
- myRaytraceSceneData.Materials.push_back (OpenGl_RaytraceMaterial());
+ myRaytraceGeometry.Materials.push_back (OpenGl_RaytraceMaterial());
}
- // Add OpenGL elements from group (only arrays of primitives)
+ // Add OpenGL elements from group (extract primitives arrays and aspects)
for (const OpenGl_ElementNode* aNode = anItg.Value()->FirstNode(); aNode != NULL; aNode = aNode->next)
{
if (TelNil == aNode->type)
if (anAspect != NULL)
{
- aMatID = static_cast<Standard_Integer> (myRaytraceSceneData.Materials.size());
+ aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size());
OpenGl_RaytraceMaterial aMaterial;
CreateMaterial (anAspect->IntFront(), aMaterial);
- myRaytraceSceneData.Materials.push_back (aMaterial);
+ myRaytraceGeometry.Materials.push_back (aMaterial);
}
}
else if (TelParray == aNode->type)
if (aPrimArray != NULL)
{
- AddRaytracePrimitiveArray (aPrimArray->PArray(), aMatID, theTransform);
+ NCollection_Handle<BVH_Object<Standard_ShortReal, 4> > aSet =
+ AddRaytracePrimitiveArray (aPrimArray->PArray(), aMatID, theTransform);
+
+ if (!aSet.IsNull())
+ myRaytraceGeometry.Objects().Append (aSet);
}
}
}
-
- anItg.Next();
}
- float* aTransform (NULL);
+ Standard_ShortReal* aTransform (NULL);
// Process all connected OpenGL structures
- OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures());
-
- while (anIts.More())
+ for (OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures()); anIts.More(); anIts.Next())
{
if (anIts.Value()->Transformation()->mat != NULL)
{
- float* aTransform = new float[16];
+ Standard_ShortReal* aTransform = new Standard_ShortReal[16];
for (Standard_Integer i = 0; i < 4; ++i)
for (Standard_Integer j = 0; j < 4; ++j)
if (anIts.Value()->IsRaytracable())
AddRaytraceStructure (anIts.Value(), aTransform != NULL ? aTransform : theTransform, theElements);
-
- anIts.Next();
}
delete[] aTransform;
// function : AddRaytracePrimitiveArray
// purpose : Adds OpenGL primitive array to ray-traced scene geometry
// =======================================================================
-Standard_Boolean OpenGl_Workspace::AddRaytracePrimitiveArray (const CALL_DEF_PARRAY* theArray,
- Standard_Integer theMatID,
- const float* theTransform)
+OpenGl_TriangleSet* OpenGl_Workspace::AddRaytracePrimitiveArray (
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theMatID, const Standard_ShortReal* theTransform)
{
if (theArray->type != TelPolygonsArrayType &&
theArray->type != TelTrianglesArrayType &&
theArray->type != TelTriangleStripsArrayType &&
theArray->type != TelQuadrangleStripsArrayType)
{
- return Standard_True;
+ return NULL;
}
if (theArray->vertices == NULL)
- return Standard_False;
+ return NULL;
#ifdef RAY_TRACE_PRINT_INFO
switch (theArray->type)
{
case TelPolygonsArrayType:
- std::cout << "\tTelPolygonsArrayType" << std::endl; break;
+ std::cout << "\tAdding TelPolygonsArrayType" << std::endl; break;
case TelTrianglesArrayType:
- std::cout << "\tTelTrianglesArrayType" << std::endl; break;
+ std::cout << "\tAdding TelTrianglesArrayType" << std::endl; break;
case TelQuadranglesArrayType:
- std::cout << "\tTelQuadranglesArrayType" << std::endl; break;
+ std::cout << "\tAdding TelQuadranglesArrayType" << std::endl; break;
case TelTriangleFansArrayType:
- std::cout << "\tTelTriangleFansArrayType" << std::endl; break;
+ std::cout << "\tAdding TelTriangleFansArrayType" << std::endl; break;
case TelTriangleStripsArrayType:
- std::cout << "\tTelTriangleStripsArrayType" << std::endl; break;
+ std::cout << "\tAdding TelTriangleStripsArrayType" << std::endl; break;
case TelQuadrangleStripsArrayType:
- std::cout << "\tTelQuadrangleStripsArrayType" << std::endl; break;
+ std::cout << "\tAdding TelQuadrangleStripsArrayType" << std::endl; break;
}
#endif
- // Simple optimization to eliminate possible memory allocations
- // during processing of the primitive array vertices
- myRaytraceSceneData.Vertices.reserve (
- myRaytraceSceneData.Vertices.size() + theArray->num_vertexs);
-
- const Standard_Integer aFirstVert = static_cast<Standard_Integer> (myRaytraceSceneData.Vertices.size());
+ OpenGl_TriangleSet* aSet = new OpenGl_TriangleSet;
- for (Standard_Integer aVert = 0; aVert < theArray->num_vertexs; ++aVert)
{
- OpenGl_RTVec4f aVertex (theArray->vertices[aVert].xyz[0],
- theArray->vertices[aVert].xyz[1],
- theArray->vertices[aVert].xyz[2],
- 1.f);
+ aSet->Vertices.reserve (theArray->num_vertexs);
- if (theTransform)
- aVertex = MatVecMult (theTransform, aVertex);
+ for (Standard_Integer aVert = 0; aVert < theArray->num_vertexs; ++aVert)
+ {
+ BVH_Vec4f aVertex (theArray->vertices[aVert].xyz[0],
+ theArray->vertices[aVert].xyz[1],
+ theArray->vertices[aVert].xyz[2],
+ 1.f);
+ if (theTransform)
+ aVertex = MatVecMult (theTransform, aVertex);
- myRaytraceSceneData.Vertices.push_back (aVertex);
+ aSet->Vertices.push_back (aVertex);
+ }
- myRaytraceSceneData.AABB.Add (aVertex);
- }
+ aSet->Normals.reserve (theArray->num_vertexs);
- myRaytraceSceneData.Normals.reserve (
- myRaytraceSceneData.Normals.size() + theArray->num_vertexs);
+ for (Standard_Integer aNorm = 0; aNorm < theArray->num_vertexs; ++aNorm)
+ {
+ BVH_Vec4f aNormal;
- for (Standard_Integer aNorm = 0; aNorm < theArray->num_vertexs; ++aNorm)
- {
- OpenGl_RTVec4f aNormal;
+ // Note: In case of absence of normals, the
+ // renderer uses generated geometric normals
- // Note: In case of absence of normals, the visualizer
- // will use generated geometric normals
+ if (theArray->vnormals != NULL)
+ {
+ aNormal = BVH_Vec4f (theArray->vnormals[aNorm].xyz[0],
+ theArray->vnormals[aNorm].xyz[1],
+ theArray->vnormals[aNorm].xyz[2],
+ 0.f);
- if (theArray->vnormals != NULL)
- {
- aNormal = OpenGl_RTVec4f (theArray->vnormals[aNorm].xyz[0],
- theArray->vnormals[aNorm].xyz[1],
- theArray->vnormals[aNorm].xyz[2],
- 0.f);
+ if (theTransform)
+ aNormal = MatVecMult (theTransform, aNormal);
+ }
- if (theTransform)
- aNormal = MatVecMult (theTransform, aNormal);
+ aSet->Normals.push_back (aNormal);
}
- myRaytraceSceneData.Normals.push_back (aNormal);
- }
+ if (theArray->num_bounds > 0)
+ {
+ #ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "\tNumber of bounds = " << theArray->num_bounds << std::endl;
+ #endif
- if (theArray->num_bounds > 0)
- {
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "\tNumber of bounds = " << theArray->num_bounds << std::endl;
-#endif
+ Standard_Integer aBoundStart = 0;
- Standard_Integer aVertOffset = 0;
+ for (Standard_Integer aBound = 0; aBound < theArray->num_bounds; ++aBound)
+ {
+ const Standard_Integer aVertNum = theArray->bounds[aBound];
- for (Standard_Integer aBound = 0; aBound < theArray->num_bounds; ++aBound)
+ #ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "\tAdding indices from bound " << aBound << ": " <<
+ aBoundStart << " .. " << aVertNum << std::endl;
+ #endif
+
+ if (!AddRaytraceVertexIndices (aSet, theArray, aBoundStart, aVertNum, theMatID))
+ {
+ delete aSet;
+ return NULL;
+ }
+
+ aBoundStart += aVertNum;
+ }
+ }
+ else
{
- const Standard_Integer aVertNum = theArray->bounds[aBound];
+ const Standard_Integer aVertNum = theArray->num_edges > 0 ? theArray->num_edges : theArray->num_vertexs;
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "\tAdd indices from bound " << aBound << ": " <<
- aVertOffset << ", " << aVertNum << std::endl;
-#endif
+ #ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "\tAdding indices from array: " << aVertNum << std::endl;
+ #endif
- if (!AddRaytraceVertexIndices (theArray, aFirstVert, aVertOffset, aVertNum, theMatID))
+ if (!AddRaytraceVertexIndices (aSet, theArray, 0, aVertNum, theMatID))
{
- return Standard_False;
+ delete aSet;
+ return NULL;
}
-
- aVertOffset += aVertNum;
}
}
- else
- {
- const Standard_Integer aVertNum = theArray->num_edges > 0 ? theArray->num_edges : theArray->num_vertexs;
-#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "\tAdd indices: " << aVertNum << std::endl;
-#endif
-
- return AddRaytraceVertexIndices (theArray, aFirstVert, 0, aVertNum, theMatID);
- }
+ if (aSet->Size() != 0)
+ aSet->MarkDirty();
- return Standard_True;
+ return aSet;
}
// =======================================================================
// function : AddRaytraceVertexIndices
// purpose : Adds vertex indices to ray-traced scene geometry
// =======================================================================
-Standard_Boolean OpenGl_Workspace::AddRaytraceVertexIndices (const CALL_DEF_PARRAY* theArray,
- Standard_Integer theFirstVert,
- Standard_Integer theVertOffset,
- Standard_Integer theVertNum,
- Standard_Integer theMatID)
+Standard_Boolean OpenGl_Workspace::AddRaytraceVertexIndices (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
{
- myRaytraceSceneData.Triangles.reserve (myRaytraceSceneData.Triangles.size() + theVertNum);
switch (theArray->type)
{
- case TelTrianglesArrayType: return AddRaytraceTriangleArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
- case TelQuadranglesArrayType: return AddRaytraceQuadrangleArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
- case TelTriangleFansArrayType: return AddRaytraceTriangleFanArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
- case TelTriangleStripsArrayType: return AddRaytraceTriangleStripArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
- case TelQuadrangleStripsArrayType: return AddRaytraceQuadrangleStripArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
- case TelPolygonsArrayType: return AddRaytracePolygonArray (theArray, theFirstVert, theVertOffset, theVertNum, theMatID);
- default: return Standard_False;
+ case TelTrianglesArrayType:
+ return AddRaytraceTriangleArray (theSet, theArray, theOffset, theCount, theMatID);
+
+ case TelQuadranglesArrayType:
+ return AddRaytraceQuadrangleArray (theSet, theArray, theOffset, theCount, theMatID);
+
+ case TelTriangleFansArrayType:
+ return AddRaytraceTriangleFanArray (theSet, theArray, theOffset, theCount, theMatID);
+
+ case TelTriangleStripsArrayType:
+ return AddRaytraceTriangleStripArray (theSet, theArray, theOffset, theCount, theMatID);
+
+ case TelQuadrangleStripsArrayType:
+ return AddRaytraceQuadrangleStripArray (theSet, theArray, theOffset, theCount, theMatID);
+
+ default:
+ return AddRaytracePolygonArray (theSet, theArray, theOffset, theCount, theMatID);
}
}
// function : AddRaytraceTriangleArray
// purpose : Adds OpenGL triangle array to ray-traced scene geometry
// =======================================================================
-Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleArray (const CALL_DEF_PARRAY* theArray,
- Standard_Integer theFirstVert,
- Standard_Integer theVertOffset,
- Standard_Integer theVertNum,
- Standard_Integer theMatID)
+Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
{
- if (theVertNum < 3)
+ if (theCount < 3)
return Standard_True;
+ theSet->Elements.reserve (theSet->Elements.size() + theCount / 3);
+
if (theArray->num_edges > 0)
{
- for (Standard_Integer aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; aVert += 3)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; aVert += 3)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[aVert + 0],
- theFirstVert + theArray->edges[aVert + 1],
- theFirstVert + theArray->edges[aVert + 2],
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
+ theArray->edges[aVert + 1],
+ theArray->edges[aVert + 2],
+ theMatID));
}
}
else
{
- for (Standard_Integer aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; aVert += 3)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; aVert += 3)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0,
- theFirstVert + aVert + 1,
- theFirstVert + aVert + 2,
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
+ aVert + 1,
+ aVert + 2,
+ theMatID));
}
}
// function : AddRaytraceTriangleFanArray
// purpose : Adds OpenGL triangle fan array to ray-traced scene geometry
// =======================================================================
-Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleFanArray (const CALL_DEF_PARRAY* theArray,
- Standard_Integer theFirstVert,
- Standard_Integer theVertOffset,
- Standard_Integer theVertNum,
- Standard_Integer theMatID)
+Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleFanArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
{
- if (theVertNum < 3)
+ if (theCount < 3)
return Standard_True;
+ theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
+
if (theArray->num_edges > 0)
{
- for (Standard_Integer aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[theVertOffset],
- theFirstVert + theArray->edges[aVert + 1],
- theFirstVert + theArray->edges[aVert + 2],
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset],
+ theArray->edges[aVert + 1],
+ theArray->edges[aVert + 2],
+ theMatID));
}
}
else
{
- for (Standard_Integer aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset,
- theFirstVert + aVert + 1,
- theFirstVert + aVert + 2,
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (theOffset,
+ aVert + 1,
+ aVert + 2,
+ theMatID));
}
}
// function : AddRaytraceTriangleStripArray
// purpose : Adds OpenGL triangle strip array to ray-traced scene geometry
// =======================================================================
-Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleStripArray (const CALL_DEF_PARRAY* theArray,
- Standard_Integer theFirstVert,
- Standard_Integer theVertOffset,
- Standard_Integer theVertNum,
- Standard_Integer theMatID)
+Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleStripArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
{
- if (theVertNum < 3)
+ if (theCount < 3)
return Standard_True;
+ theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
+
if (theArray->num_edges > 0)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
- theFirstVert + theArray->edges[theVertOffset + 0],
- theFirstVert + theArray->edges[theVertOffset + 1],
- theFirstVert + theArray->edges[theVertOffset + 2],
- theMatID));
-
- for (Standard_Integer aVert = theVertOffset + 1, aTriNum = 1; aVert < theVertOffset + theVertNum - 2; ++aVert, ++aTriNum)
+ for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
- theFirstVert + theArray->edges[aVert + (aTriNum % 2) ? 1 : 0],
- theFirstVert + theArray->edges[aVert + (aTriNum % 2) ? 0 : 1],
- theFirstVert + theArray->edges[aVert + 2],
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + aCW ? 1 : 0],
+ theArray->edges[aVert + aCW ? 0 : 1],
+ theArray->edges[aVert + 2],
+ theMatID));
}
}
else
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset + 0,
- theFirstVert + theVertOffset + 1,
- theFirstVert + theVertOffset + 2,
- theMatID));
-
- for (Standard_Integer aVert = theVertOffset + 1, aTriNum = 1; aVert < theVertOffset + theVertNum - 2; ++aVert, ++aTriNum)
+ for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + ( aTriNum % 2 ) ? 1 : 0,
- theFirstVert + aVert + ( aTriNum % 2 ) ? 0 : 1,
- theFirstVert + aVert + 2,
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (aVert + aCW ? 1 : 0,
+ aVert + aCW ? 0 : 1,
+ aVert + 2,
+ theMatID));
}
}
// function : AddRaytraceQuadrangleArray
// purpose : Adds OpenGL quad array to ray-traced scene geometry
// =======================================================================
-Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleArray (const CALL_DEF_PARRAY* theArray,
- Standard_Integer theFirstVert,
- Standard_Integer theVertOffset,
- Standard_Integer theVertNum,
- Standard_Integer theMatID)
+Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
{
- if (theVertNum < 4)
+ if (theCount < 4)
return Standard_True;
+ theSet->Elements.reserve (theSet->Elements.size() + theCount / 2);
+
if (theArray->num_edges > 0)
{
- for (Standard_Integer aVert = theVertOffset; aVert < theVertOffset + theVertNum - 3; aVert += 4)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 4)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[aVert + 0],
- theFirstVert + theArray->edges[aVert + 1],
- theFirstVert + theArray->edges[aVert + 2],
- theMatID));
-
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[aVert + 0],
- theFirstVert + theArray->edges[aVert + 2],
- theFirstVert + theArray->edges[aVert + 3],
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
+ theArray->edges[aVert + 1],
+ theArray->edges[aVert + 2],
+ theMatID));
+
+ theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
+ theArray->edges[aVert + 2],
+ theArray->edges[aVert + 3],
+ theMatID));
}
}
else
{
- for (Standard_Integer aVert = theVertOffset; aVert < theVertOffset + theVertNum - 3; aVert += 4)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 4)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0,
- theFirstVert + aVert + 1,
- theFirstVert + aVert + 2,
- theMatID));
-
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0,
- theFirstVert + aVert + 2,
- theFirstVert + aVert + 3,
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
+ aVert + 1,
+ aVert + 2,
+ theMatID));
+
+ theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
+ aVert + 2,
+ aVert + 3,
+ theMatID));
}
}
// function : AddRaytraceQuadrangleStripArray
// purpose : Adds OpenGL quad strip array to ray-traced scene geometry
// =======================================================================
-Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleStripArray (const CALL_DEF_PARRAY* theArray,
- Standard_Integer theFirstVert,
- Standard_Integer theVertOffset,
- Standard_Integer theVertNum,
- Standard_Integer theMatID)
+Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleStripArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
{
- if (theVertNum < 4)
+ if (theCount < 4)
return Standard_True;
+ theSet->Elements.reserve (theSet->Elements.size() + 2 * theCount - 6);
+
if (theArray->num_edges > 0)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
- theFirstVert + theArray->edges[theVertOffset + 0],
- theFirstVert + theArray->edges[theVertOffset + 1],
- theFirstVert + theArray->edges[theVertOffset + 2],
- theMatID));
-
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
- theFirstVert + theArray->edges[theVertOffset + 1],
- theFirstVert + theArray->edges[theVertOffset + 3],
- theFirstVert + theArray->edges[theVertOffset + 2],
- theMatID));
-
- for (Standard_Integer aVert = theVertOffset + 2; aVert < theVertOffset + theVertNum - 3; aVert += 2)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 2)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
- theFirstVert + theArray->edges[aVert + 0],
- theFirstVert + theArray->edges[aVert + 1],
- theFirstVert + theArray->edges[aVert + 2],
- theMatID));
-
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (
- theFirstVert + theArray->edges[aVert + 1],
- theFirstVert + theArray->edges[aVert + 3],
- theFirstVert + theArray->edges[aVert + 2],
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0],
+ theArray->edges[aVert + 1],
+ theArray->edges[aVert + 2],
+ theMatID));
+
+ theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 1],
+ theArray->edges[aVert + 3],
+ theArray->edges[aVert + 2],
+ theMatID));
}
}
else
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + 0,
- theFirstVert + 1,
- theFirstVert + 2,
- theMatID));
-
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + 1,
- theFirstVert + 3,
- theFirstVert + 2,
- theMatID));
-
- for (Standard_Integer aVert = theVertOffset + 2; aVert < theVertOffset + theVertNum - 3; aVert += 2)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 2)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 0,
- theFirstVert + aVert + 1,
- theFirstVert + aVert + 2,
- theMatID));
-
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + aVert + 1,
- theFirstVert + aVert + 3,
- theFirstVert + aVert + 2,
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (aVert + 0,
+ aVert + 1,
+ aVert + 2,
+ theMatID));
+
+ theSet->Elements.push_back (BVH_Vec4i (aVert + 1,
+ aVert + 3,
+ aVert + 2,
+ theMatID));
}
}
// function : AddRaytracePolygonArray
// purpose : Adds OpenGL polygon array to ray-traced scene geometry
// =======================================================================
-Standard_Boolean OpenGl_Workspace::AddRaytracePolygonArray (const CALL_DEF_PARRAY* theArray,
- Standard_Integer theFirstVert,
- Standard_Integer theVertOffset,
- Standard_Integer theVertNum,
- Standard_Integer theMatID)
+Standard_Boolean OpenGl_Workspace::AddRaytracePolygonArray (OpenGl_TriangleSet* theSet,
+ const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID)
{
- if (theArray->num_vertexs < 3)
+ if (theCount < 3)
return Standard_True;
- if (theArray->edges != NULL)
+ theSet->Elements.reserve (theSet->Elements.size() + theCount - 2);
+
+ if (theArray->num_edges > 0)
{
- for (Standard_Integer aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theArray->edges[theVertOffset],
- theFirstVert + theArray->edges[aVert + 1],
- theFirstVert + theArray->edges[aVert + 2],
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset],
+ theArray->edges[aVert + 1],
+ theArray->edges[aVert + 2],
+ theMatID));
}
}
else
{
- for (Standard_Integer aVert = theVertOffset; aVert < theVertOffset + theVertNum - 2; ++aVert)
+ for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert)
{
- myRaytraceSceneData.Triangles.push_back (OpenGl_RTVec4i (theFirstVert + theVertOffset,
- theFirstVert + aVert + 1,
- theFirstVert + aVert + 2,
- theMatID));
+ theSet->Elements.push_back (BVH_Vec4i (theOffset,
+ aVert + 1,
+ aVert + 2,
+ theMatID));
}
}
// =======================================================================
Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble theInvModelView[16])
{
- myRaytraceSceneData.LightSources.clear();
+ myRaytraceGeometry.Sources.clear();
- OpenGl_RTVec4f anAmbient (0.0f, 0.0f, 0.0f, 0.0f);
- for (OpenGl_ListOfLight::Iterator anItl (myView->LightList());
- anItl.More(); anItl.Next())
+ myRaytraceGeometry.GlobalAmbient = BVH_Vec4f (0.0f, 0.0f, 0.0f, 0.0f);
+
+ for (OpenGl_ListOfLight::Iterator anItl (myView->LightList()); anItl.More(); anItl.Next())
{
const OpenGl_Light& aLight = anItl.Value();
+
if (aLight.Type == Visual3d_TOLS_AMBIENT)
{
- anAmbient += OpenGl_RTVec4f (aLight.Color.r(), aLight.Color.g(), aLight.Color.b(), 0.0f);
+ myRaytraceGeometry.GlobalAmbient += BVH_Vec4f (aLight.Color.r(),
+ aLight.Color.g(),
+ aLight.Color.b(),
+ 0.0f);
continue;
}
- OpenGl_RTVec4f aDiffuse (aLight.Color.r(), aLight.Color.g(), aLight.Color.b(), 1.0f);
- OpenGl_RTVec4f aPosition (-aLight.Direction.x(), -aLight.Direction.y(), -aLight.Direction.z(), 0.0f);
+ BVH_Vec4f aDiffuse (aLight.Color.r(),
+ aLight.Color.g(),
+ aLight.Color.b(),
+ 1.0f);
+
+ BVH_Vec4f aPosition (-aLight.Direction.x(),
+ -aLight.Direction.y(),
+ -aLight.Direction.z(),
+ 0.0f);
+
if (aLight.Type != Visual3d_TOLS_DIRECTIONAL)
{
- aPosition = OpenGl_RTVec4f (aLight.Position.x(), aLight.Position.y(), aLight.Position.z(), 1.0f);
+ aPosition = BVH_Vec4f (aLight.Position.x(),
+ aLight.Position.y(),
+ aLight.Position.z(),
+ 1.0f);
}
+
if (aLight.IsHeadlight)
- {
aPosition = MatVecMult (theInvModelView, aPosition);
- }
- myRaytraceSceneData.LightSources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition));
- }
-
- if (myRaytraceSceneData.LightSources.size() > 0)
- {
- myRaytraceSceneData.LightSources.front().Ambient += anAmbient;
- }
- else
- {
- myRaytraceSceneData.LightSources.push_back (OpenGl_RaytraceLight (OpenGl_RTVec4f (anAmbient.rgb(), -1.0f)));
+ myRaytraceGeometry.Sources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition));
}
cl_int anError = CL_SUCCESS;
if (myRaytraceLightSourceBuffer != NULL)
clReleaseMemObject (myRaytraceLightSourceBuffer);
- const size_t myLightBufferSize = myRaytraceSceneData.LightSources.size() > 0
- ? myRaytraceSceneData.LightSources.size()
- : 1;
+ Standard_Integer aLightBufferSize = myRaytraceGeometry.Sources.size() != 0 ?
+ static_cast<Standard_Integer> (myRaytraceGeometry.Sources.size()) : 1;
myRaytraceLightSourceBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
- myLightBufferSize * sizeof(OpenGl_RaytraceLight),
- NULL, &anError);
+ aLightBufferSize * sizeof(OpenGl_RaytraceLight), NULL, &anError);
- if (myRaytraceSceneData.LightSources.size() > 0)
+ if (myRaytraceGeometry.Sources.size() != 0)
{
- const void* aDataPtr = myRaytraceSceneData.LightSources.front().Packed();
- anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0,
- myLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr,
- 0, NULL, NULL);
+ const void* aDataPtr = myRaytraceGeometry.Sources.front().Packed();
+
+ anError |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0,
+ aLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr, 0, NULL, NULL);
}
#ifdef RAY_TRACE_PRINT_INFO
if (anError != CL_SUCCESS)
{
- std::cout << "Error! Failed to set light sources!";
+ std::cout << "Error! Failed to set light sources";
return Standard_False;
}
}
// Create OpenCL ray tracing kernels
- myRaytraceRenderKernel = clCreateKernel (myRaytraceProgram, "Main", &anError);
+ myRaytraceRenderKernel = clCreateKernel (myRaytraceProgram, "RaytraceRender", &anError);
if (anError != CL_SUCCESS)
{
myComputeInitStatus = OpenGl_CLIS_FAIL;
return Standard_False;
}
- myRaytraceSmoothKernel = clCreateKernel (myRaytraceProgram, "MainAntialiased", &anError);
+ myRaytraceSmoothKernel = clCreateKernel (myRaytraceProgram, "RaytraceSmooth", &anError);
if (anError != CL_SUCCESS)
{
myComputeInitStatus = OpenGl_CLIS_FAIL;
// Note: For profiling set CL_QUEUE_PROFILING_ENABLE
cl_command_queue_properties aProps = CL_QUEUE_PROFILING_ENABLE;
- myRaytraceQueue = clCreateCommandQueue (myComputeContext, aDevice, aProps, &anError);
+ myComputeQueue = clCreateCommandQueue (myComputeContext, aDevice, aProps, &anError);
if (anError != CL_SUCCESS)
{
myComputeInitStatus = OpenGl_CLIS_FAIL;
clReleaseKernel (myRaytraceSmoothKernel);
clReleaseProgram (myRaytraceProgram);
- clReleaseCommandQueue (myRaytraceQueue);
-
+ clReleaseCommandQueue (myComputeQueue);
+
clReleaseMemObject (myRaytraceOutputImage);
clReleaseMemObject (myRaytraceEnvironment);
clReleaseMemObject (myRaytraceOutputImageAA);
- clReleaseMemObject (myRaytraceVertexBuffer);
- clReleaseMemObject (myRaytraceNormalBuffer);
- clReleaseMemObject (myRaytraceTriangleBuffer);
-
clReleaseMemObject (myRaytraceMaterialBuffer);
clReleaseMemObject (myRaytraceLightSourceBuffer);
- clReleaseMemObject (myRaytraceNodeMinPointBuffer);
- clReleaseMemObject (myRaytraceNodeMaxPointBuffer);
- clReleaseMemObject (myRaytraceNodeDataRcrdBuffer);
+ clReleaseMemObject (mySceneNodeInfoBuffer);
+ clReleaseMemObject (mySceneMinPointBuffer);
+ clReleaseMemObject (mySceneMaxPointBuffer);
+
+ clReleaseMemObject (myObjectNodeInfoBuffer);
+ clReleaseMemObject (myObjectMinPointBuffer);
+ clReleaseMemObject (myObjectMaxPointBuffer);
+
+ clReleaseMemObject (myGeometryVertexBuffer);
+ clReleaseMemObject (myGeometryNormalBuffer);
+ clReleaseMemObject (myGeometryTriangBuffer);
clReleaseContext (myComputeContext);
// =======================================================================
// function : WriteRaytraceSceneToDevice
-// purpose : Writes scene geometry to OpenCl device
+// purpose : Writes scene geometry to OpenCL device
// =======================================================================
Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice()
{
if (myComputeContext == NULL)
return Standard_False;
- cl_int anError = CL_SUCCESS;
+ cl_int anErrorRes = CL_SUCCESS;
+
+ if (mySceneNodeInfoBuffer != NULL)
+ anErrorRes |= clReleaseMemObject (mySceneNodeInfoBuffer);
+
+ if (mySceneMinPointBuffer != NULL)
+ anErrorRes |= clReleaseMemObject (mySceneMinPointBuffer);
+
+ if (mySceneMaxPointBuffer != NULL)
+ anErrorRes |= clReleaseMemObject (mySceneMaxPointBuffer);
- if (myRaytraceNormalBuffer != NULL)
- anError |= clReleaseMemObject (myRaytraceNormalBuffer);
+ if (myObjectNodeInfoBuffer != NULL)
+ anErrorRes |= clReleaseMemObject (myObjectNodeInfoBuffer);
- if (myRaytraceVertexBuffer != NULL)
- anError |= clReleaseMemObject (myRaytraceVertexBuffer);
+ if (myObjectMinPointBuffer != NULL)
+ anErrorRes |= clReleaseMemObject (myObjectMinPointBuffer);
- if (myRaytraceTriangleBuffer != NULL)
- anError |= clReleaseMemObject (myRaytraceTriangleBuffer);
+ if (myObjectMaxPointBuffer != NULL)
+ anErrorRes |= clReleaseMemObject (myObjectMaxPointBuffer);
- if (myRaytraceNodeMinPointBuffer != NULL)
- anError |= clReleaseMemObject (myRaytraceNodeMinPointBuffer);
+ if (myGeometryVertexBuffer != NULL)
+ anErrorRes |= clReleaseMemObject (myGeometryVertexBuffer);
- if (myRaytraceNodeMaxPointBuffer != NULL)
- anError |= clReleaseMemObject (myRaytraceNodeMaxPointBuffer);
+ if (myGeometryNormalBuffer != NULL)
+ anErrorRes |= clReleaseMemObject (myGeometryNormalBuffer);
- if (myRaytraceNodeDataRcrdBuffer != NULL)
- anError |= clReleaseMemObject (myRaytraceNodeDataRcrdBuffer);
+ if (myGeometryTriangBuffer != NULL)
+ anErrorRes |= clReleaseMemObject (myGeometryTriangBuffer);
if (myRaytraceMaterialBuffer != NULL)
- anError |= clReleaseMemObject (myRaytraceMaterialBuffer);
+ anErrorRes |= clReleaseMemObject (myRaytraceMaterialBuffer);
- if (anError != CL_SUCCESS)
+ if (anErrorRes != CL_SUCCESS)
{
#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to release OpenCL scene buffers!" << std::endl;
+ std::cout << "Error! Failed to release OpenCL buffers" << std::endl;
#endif
return Standard_False;
}
- // Create geometry buffers
- cl_int anErrorTemp = CL_SUCCESS;
- const size_t myVertexBufferSize = myRaytraceSceneData.Vertices.size() > 0
- ? myRaytraceSceneData.Vertices.size() : 1;
-
- myRaytraceVertexBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
- myVertexBufferSize * sizeof(cl_float4), NULL, &anErrorTemp);
- anError |= anErrorTemp;
-
- const size_t myNormalBufferSize = myRaytraceSceneData.Normals.size() > 0
- ? myRaytraceSceneData.Normals.size() : 1;
- myRaytraceNormalBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
- myNormalBufferSize * sizeof(cl_float4), NULL, &anErrorTemp);
- anError |= anErrorTemp;
-
- const size_t myTriangleBufferSize = myRaytraceSceneData.Triangles.size() > 0
- ? myRaytraceSceneData.Triangles.size() : 1;
- myRaytraceTriangleBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
- myTriangleBufferSize * sizeof(cl_int4), NULL, &anErrorTemp);
- anError |= anErrorTemp;
- if (anError != CL_SUCCESS)
+ /////////////////////////////////////////////////////////////////////////////
+ // Create material buffer
+
+ const size_t aMaterialBufferSize =
+ myRaytraceGeometry.Materials.size() != 0 ? myRaytraceGeometry.Materials.size() : 1;
+
+ myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL, &anErrorRes);
+
+ if (anErrorRes != CL_SUCCESS)
{
#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to create OpenCL geometry buffers!" << std::endl;
+ std::cout << "Error! Failed to create OpenCL material buffer" << std::endl;
#endif
return Standard_False;
}
- // Create material buffer
- const size_t myMaterialBufferSize = myRaytraceSceneData.Materials.size() > 0
- ? myRaytraceSceneData.Materials.size() : 1;
- myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
- myMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL,
- &anErrorTemp);
- if (anErrorTemp != CL_SUCCESS)
+ /////////////////////////////////////////////////////////////////////////////
+ // Create BVHs buffers
+
+ cl_int anErrorTmp = CL_SUCCESS;
+
+ const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = myRaytraceGeometry.BVH();
+
+ const size_t aSceneMinPointBufferSize =
+ aBVH->MinPointBuffer().size() != 0 ? aBVH->MinPointBuffer().size() : 1;
+
+ mySceneMinPointBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aSceneMinPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp);
+ anErrorRes |= anErrorTmp;
+
+ const size_t aSceneMaxPointBufferSize =
+ aBVH->MaxPointBuffer().size() != 0 ? aBVH->MaxPointBuffer().size() : 1;
+
+ mySceneMaxPointBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aSceneMaxPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp);
+ anErrorRes |= anErrorTmp;
+
+ const size_t aSceneNodeInfoBufferSize =
+ aBVH->NodeInfoBuffer().size() != 0 ? aBVH->NodeInfoBuffer().size() : 1;
+
+ mySceneNodeInfoBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aSceneNodeInfoBufferSize * sizeof(cl_int4), NULL, &anErrorTmp);
+ anErrorRes |= anErrorTmp;
+
+ if (anErrorRes != CL_SUCCESS)
{
#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to create OpenCL material buffer!" << std::endl;
+ std::cout << "Error! Failed to create OpenCL buffers for high-level scene BVH" << std::endl;
#endif
return Standard_False;
}
- // Create BVH buffers
- OpenGl_BVH aTree = myBVHBuilder.Tree();
- const size_t myNodeMinPointBufferSize = aTree.MinPointBuffer().size() > 0
- ? aTree.MinPointBuffer().size() : 1;
- myRaytraceNodeMinPointBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
- myNodeMinPointBufferSize * sizeof(cl_float4), NULL,
- &anErrorTemp);
- anError |= anErrorTemp;
-
- const size_t myNodeMaxPointBufferSize = aTree.MaxPointBuffer().size() > 0
- ? aTree.MaxPointBuffer().size() : 1;
- myRaytraceNodeMaxPointBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
- myNodeMaxPointBufferSize * sizeof(cl_float4), NULL,
- &anError);
- anError |= anErrorTemp;
-
- const size_t myNodeDataRecordBufferSize = aTree.DataRcrdBuffer().size() > 0
- ? aTree.DataRcrdBuffer().size() : 1;
- myRaytraceNodeDataRcrdBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY,
- myNodeDataRecordBufferSize * sizeof(cl_int4), NULL,
- &anError);
- anError |= anErrorTemp;
- if (anError != CL_SUCCESS)
+ Standard_Integer aTotalVerticesNb = 0;
+ Standard_Integer aTotalElementsNb = 0;
+ Standard_Integer aTotalBVHNodesNb = 0;
+
+ for (Standard_Integer anElemIndex = 0; anElemIndex < myRaytraceGeometry.Size(); ++anElemIndex)
+ {
+ OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
+ myRaytraceGeometry.Objects().ChangeValue (anElemIndex).operator->());
+
+ Standard_ASSERT_RETURN (aTriangleSet != NULL,
+ "Error! Failed to get triangulation of OpenGL element", Standard_False);
+
+ aTotalVerticesNb += aTriangleSet->Vertices.size();
+ aTotalElementsNb += aTriangleSet->Elements.size();
+
+ Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(),
+ "Error! Failed to get bottom-level BVH of OpenGL element", Standard_False);
+
+ aTotalBVHNodesNb += aTriangleSet->BVH()->NodeInfoBuffer().size();
+ }
+
+ aTotalBVHNodesNb = aTotalBVHNodesNb > 0 ? aTotalBVHNodesNb : 1;
+
+ myObjectNodeInfoBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_int4), NULL, &anErrorTmp);
+ anErrorRes |= anErrorTmp;
+
+ myObjectMinPointBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp);
+ anErrorRes |= anErrorTmp;
+
+ myObjectMaxPointBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp);
+ anErrorRes |= anErrorTmp;
+
+ if (anErrorRes != CL_SUCCESS)
{
#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to create OpenCL BVH buffers!" << std::endl;
+ std::cout << "Error! Failed to create OpenCL buffers for bottom-level scene BVHs" << std::endl;
#endif
return Standard_False;
}
- // Write scene geometry buffers
- if (myRaytraceSceneData.Triangles.size() > 0)
- {
- anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceVertexBuffer, CL_FALSE,
- 0, myRaytraceSceneData.Vertices.size() * sizeof(cl_float4),
- &myRaytraceSceneData.Vertices.front(),
- 0, NULL, NULL);
- anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNormalBuffer, CL_FALSE,
- 0, myRaytraceSceneData.Normals.size() * sizeof(cl_float4),
- &myRaytraceSceneData.Normals.front(),
- 0, NULL, NULL);
- anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceTriangleBuffer, CL_FALSE,
- 0, myRaytraceSceneData.Triangles.size() * sizeof(cl_int4),
- &myRaytraceSceneData.Triangles.front(),
- 0, NULL, NULL);
- if (anError != CL_SUCCESS)
+ /////////////////////////////////////////////////////////////////////////////
+ // Create geometry buffers
+
+ aTotalVerticesNb = aTotalVerticesNb > 0 ? aTotalVerticesNb : 1;
+
+ myGeometryVertexBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp);
+ anErrorRes |= anErrorTmp;
+
+ myGeometryNormalBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp);
+ anErrorRes |= anErrorTmp;
+
+ aTotalElementsNb = aTotalElementsNb > 0 ? aTotalElementsNb : 1;
+
+ myGeometryTriangBuffer = clCreateBuffer (myComputeContext,
+ CL_MEM_READ_ONLY, aTotalElementsNb * sizeof(cl_int4), NULL, &anErrorTmp);
+ anErrorRes |= anErrorTmp;
+
+ if (anErrorRes != CL_SUCCESS)
+ {
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error! Failed to create OpenCL geometry buffers" << std::endl;
+#endif
+ return Standard_False;
+ }
+
+ /////////////////////////////////////////////////////////////////////////////
+ // Write BVH and geometry buffers
+
+ if (aBVH->NodeInfoBuffer().size() != 0)
+ {
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneNodeInfoBuffer, CL_FALSE, 0,
+ aSceneNodeInfoBufferSize * sizeof(cl_int4), &aBVH->NodeInfoBuffer().front(), 0, NULL, NULL);
+
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMinPointBuffer, CL_FALSE, 0,
+ aSceneMinPointBufferSize * sizeof(cl_float4), &aBVH->MinPointBuffer().front(), 0, NULL, NULL);
+
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMaxPointBuffer, CL_FALSE, 0,
+ aSceneMaxPointBufferSize * sizeof(cl_float4), &aBVH->MaxPointBuffer().front(), 0, NULL, NULL);
+
+ anErrorRes |= clFinish (myComputeQueue);
+
+ if (anErrorRes != CL_SUCCESS)
{
- #ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to write OpenCL geometry buffers!" << std::endl;
- #endif
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error! Failed to write OpenCL buffers for high-level scene BVH" << std::endl;
+#endif
return Standard_False;
}
- }
- // Write BVH buffers
- if (aTree.DataRcrdBuffer().size() > 0)
- {
- anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeMinPointBuffer, CL_FALSE,
- 0, aTree.MinPointBuffer().size() * sizeof(cl_float4),
- &aTree.MinPointBuffer().front(),
- 0, NULL, NULL);
- anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeMaxPointBuffer, CL_FALSE,
- 0, aTree.MaxPointBuffer().size() * sizeof(cl_float4),
- &aTree.MaxPointBuffer().front(),
- 0, NULL, NULL);
- anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceNodeDataRcrdBuffer, CL_FALSE,
- 0, aTree.DataRcrdBuffer().size() * sizeof(cl_int4),
- &aTree.DataRcrdBuffer().front(),
- 0, NULL, NULL);
- if (anError != CL_SUCCESS)
+ for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx)
{
- #ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to write OpenCL BVH buffers!" << std::endl;
- #endif
- return Standard_False;
+ if (!aBVH->IsOuter (aNodeIdx))
+ continue;
+
+ OpenGl_TriangleSet* aTriangleSet = myRaytraceGeometry.TriangleSet (aNodeIdx);
+
+ Standard_ASSERT_RETURN (aTriangleSet != NULL,
+ "Error! Failed to get triangulation of OpenGL element", Standard_False);
+
+ const size_t aBVHBuffserSize =
+ aTriangleSet->BVH()->NodeInfoBuffer().size() != 0 ? aTriangleSet->BVH()->NodeInfoBuffer().size() : 1;
+
+ const Standard_Integer aBVHOffset = myRaytraceGeometry.AccelerationOffset (aNodeIdx);
+
+ Standard_ASSERT_RETURN (aBVHOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
+ "Error! Failed to get offset for bottom-level BVH", Standard_False);
+
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectNodeInfoBuffer, CL_FALSE, aBVHOffset * sizeof(cl_int4),
+ aBVHBuffserSize * sizeof(cl_int4), &aTriangleSet->BVH()->NodeInfoBuffer().front(), 0, NULL, NULL);
+
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMinPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4),
+ aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MinPointBuffer().front(), 0, NULL, NULL);
+
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMaxPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4),
+ aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MaxPointBuffer().front(), 0, NULL, NULL);
+
+ anErrorRes |= clFinish (myComputeQueue);
+
+ if (anErrorRes != CL_SUCCESS)
+ {
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error! Failed to write OpenCL buffers for bottom-level scene BVHs" << std::endl;
+#endif
+ return Standard_False;
+ }
+
+ const Standard_Integer aVerticesOffset = myRaytraceGeometry.VerticesOffset (aNodeIdx);
+
+ Standard_ASSERT_RETURN (aVerticesOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
+ "Error! Failed to get offset for triangulation vertices of OpenGL element", Standard_False);
+
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryVertexBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4),
+ aTriangleSet->Vertices.size() * sizeof(cl_float4), &aTriangleSet->Vertices.front(), 0, NULL, NULL);
+
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryNormalBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4),
+ aTriangleSet->Normals.size() * sizeof(cl_float4), &aTriangleSet->Normals.front(), 0, NULL, NULL);
+
+ const Standard_Integer anElementsOffset = myRaytraceGeometry.ElementsOffset (aNodeIdx);
+
+ Standard_ASSERT_RETURN (anElementsOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET,
+ "Error! Failed to get offset for triangulation elements of OpenGL element", Standard_False);
+
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryTriangBuffer, CL_FALSE, anElementsOffset * sizeof(cl_int4),
+ aTriangleSet->Elements.size() * sizeof(cl_int4), &aTriangleSet->Elements.front(), 0, NULL, NULL);
+
+ anErrorRes |= clFinish (myComputeQueue);
+
+ if (anErrorRes != CL_SUCCESS)
+ {
+#ifdef RAY_TRACE_PRINT_INFO
+ std::cout << "Error! Failed to write OpenCL triangulation buffers for OpenGL element" << std::endl;
+#endif
+ return Standard_False;
+ }
}
}
- // Write material buffers
- if (myRaytraceSceneData.Materials.size() > 0)
+ /////////////////////////////////////////////////////////////////////////////
+ // Write material buffer
+
+ if (myRaytraceGeometry.Materials.size() != 0)
{
- const size_t aSize = myRaytraceSceneData.Materials.size();
- const void* aDataPtr = myRaytraceSceneData.Materials.front().Packed();
+ const void* aDataPtr = myRaytraceGeometry.Materials.front().Packed();
- anError |= clEnqueueWriteBuffer (myRaytraceQueue, myRaytraceMaterialBuffer, CL_FALSE,
- 0, aSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr,
- 0, NULL, NULL);
- if (anError != CL_SUCCESS)
+ anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceMaterialBuffer,
+ CL_FALSE, 0, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr, 0, NULL, NULL);
+
+ if (anErrorRes != CL_SUCCESS)
{
#ifdef RAY_TRACE_PRINT_INFO
- std::cout << "Error! Failed to write OpenCL material buffer!" << std::endl;
+ std::cout << "Error! Failed to write OpenCL material buffer" << std::endl;
#endif
return Standard_False;
}
}
- anError |= clFinish (myRaytraceQueue);
+ anErrorRes |= clFinish (myComputeQueue);
+
+ if (anErrorRes == CL_SUCCESS)
+ {
+ myIsRaytraceDataValid = myRaytraceGeometry.Objects().Size() != 0;
+ }
#ifdef RAY_TRACE_PRINT_INFO
- if (anError != CL_SUCCESS)
- std::cout << "Error! Failed to set scene data buffers!" << std::endl;
+ else
+ {
+ std::cout << "Error! Failed to set scene data buffers" << std::endl;
+ }
#endif
- if (anError == CL_SUCCESS)
- myIsRaytraceDataValid = myRaytraceSceneData.Triangles.size() > 0;
-
#ifdef RAY_TRACE_PRINT_INFO
+
+ Standard_ShortReal aMemUsed = 0.f;
+
+ for (Standard_Integer anElemIdx = 0; anElemIdx < myRaytraceGeometry.Size(); ++anElemIdx)
+ {
+ OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> (
+ myRaytraceGeometry.Objects().ChangeValue (anElemIdx).operator->());
- float aMemUsed = static_cast<float> (
- myRaytraceSceneData.Materials.size()) * sizeof (OpenGl_RaytraceMaterial);
-
- aMemUsed += static_cast<float> (
- myRaytraceSceneData.Triangles.size() * sizeof (OpenGl_RTVec4i) +
- myRaytraceSceneData.Vertices.size() * sizeof (OpenGl_RTVec4f) +
- myRaytraceSceneData.Normals.size() * sizeof (OpenGl_RTVec4f));
+ 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<float> (
- aTree.MinPointBuffer().size() * sizeof (OpenGl_RTVec4f) +
- aTree.MaxPointBuffer().size() * sizeof (OpenGl_RTVec4f) +
- aTree.DataRcrdBuffer().size() * sizeof (OpenGl_RTVec4i));
+ 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));
+ }
+
+ 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 / 1e6f << std::endl;
+ std::cout << "GPU Memory Used (MB): ~" << aMemUsed / 1048576 << std::endl;
#endif
- myRaytraceSceneData.Clear();
-
- myBVHBuilder.CleanUp();
-
- return (CL_SUCCESS == anError);
+ return (CL_SUCCESS == anErrorRes);
}
-#define OPENCL_GROUP_SIZE_TEST_
+// Use it to estimate the optimal size of OpenCL work group
+// #define OPENCL_GROUP_SIZE_TEST
// =======================================================================
// function : RunRaytraceOpenCLKernels
// purpose : Runs OpenCL ray-tracing kernels
// =======================================================================
-Standard_Boolean OpenGl_Workspace::RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView,
- const GLfloat theOrigins[16],
- const GLfloat theDirects[16],
- const Standard_Integer theSizeX,
- const Standard_Integer theSizeY)
+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 || myRaytraceQueue == NULL)
+ if (myRaytraceRenderKernel == NULL || myComputeQueue == NULL)
return Standard_False;
- ////////////////////////////////////////////////////////////
+ ////////////////////////////////////////////////////////////////////////
// Set kernel arguments
cl_uint anIndex = 0;
cl_int anError = 0;
- anError = clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceOutputImage);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceEnvironment);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceNodeMinPointBuffer);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceNodeMaxPointBuffer);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceNodeDataRcrdBuffer);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceLightSourceBuffer);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceMaterialBuffer);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceVertexBuffer);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceNormalBuffer);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceTriangleBuffer);
-
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_float16), theOrigins);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_float16), theDirects);
-
- cl_int aLightCount = static_cast<cl_int> (myRaytraceSceneData.LightSources.size());
-
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_int), &aLightCount);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_float), &myRaytraceSceneEpsilon);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_float), &myRaytraceSceneRadius);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_int), &theCView.IsShadowsEnabled);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_int), &theCView.IsReflectionsEnabled);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_int), &theSizeX);
- anError |= clSetKernelArg (myRaytraceRenderKernel, anIndex++,
- sizeof(cl_int), &theSizeY);
+ cl_int aLightSourceBufferSize = 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 aMsg = "Error! Failed to set arguments of ray-tracing kernel!";
+ const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of ray-tracing kernel!";
+
myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
- GL_DEBUG_TYPE_ERROR_ARB,
- 0,
- GL_DEBUG_SEVERITY_HIGH_ARB,
- aMsg);
+ GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+
return Standard_False;
}
- // Note: second-pass 'smoothing' kernel runs only if anti-aliasing is enabled
+ // Second-pass 'smoothing' kernel runs only if anti-aliasing is enabled
if (theCView.IsAntialiasingEnabled)
{
anIndex = 0;
- anError = clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceOutputImage);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceOutputImageAA);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceEnvironment);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceNodeMinPointBuffer);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceNodeMaxPointBuffer);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceNodeDataRcrdBuffer);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceLightSourceBuffer);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceMaterialBuffer);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceVertexBuffer);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceNormalBuffer);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_mem), &myRaytraceTriangleBuffer);
-
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_float16), theOrigins);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_float16), theDirects);
-
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_int), &aLightCount);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_float), &myRaytraceSceneEpsilon);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_float), &myRaytraceSceneRadius);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_int), &theCView.IsShadowsEnabled);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_int), &theCView.IsReflectionsEnabled);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_int), &theSizeX);
- anError |= clSetKernelArg (myRaytraceSmoothKernel, anIndex++,
- sizeof(cl_int), &theSizeY);
+
+ 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 aMsg = "Error! Failed to set arguments of 'smoothing' kernel!";
+ const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of smoothing kernel!";
+
myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
- GL_DEBUG_TYPE_ERROR_ARB,
- 0,
- GL_DEBUG_SEVERITY_HIGH_ARB,
- aMsg);
+ GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+
return Standard_False;
}
}
+ ////////////////////////////////////////////////////////////////////////
// Set work size
- size_t aLocSizeRender[] = { myIsAmdComputePlatform ? 2 : 4, 32 };
+
+ size_t aLocWorkSize[] = { myIsAmdComputePlatform ? 2 : 4, 32 };
#ifdef OPENCL_GROUP_SIZE_TEST
- for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1 )
- for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1 )
-#endif
+ for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1)
+ for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1)
{
-#ifdef OPENCL_GROUP_SIZE_TEST
- aLocSizeRender[0] = aLocX;
- aLocSizeRender[1] = aLocY;
+ aLocWorkSize[0] = aLocX;
+ aLocWorkSize[1] = aLocY;
#endif
size_t aWorkSizeX = theSizeX;
- if (aWorkSizeX % aLocSizeRender[0] != 0)
- aWorkSizeX += aLocSizeRender[0] - aWorkSizeX % aLocSizeRender[0];
+ if (aWorkSizeX % aLocWorkSize[0] != 0)
+ aWorkSizeX += aLocWorkSize[0] - aWorkSizeX % aLocWorkSize[0];
size_t aWokrSizeY = theSizeY;
- if (aWokrSizeY % aLocSizeRender[1] != 0 )
- aWokrSizeY += aLocSizeRender[1] - aWokrSizeY % aLocSizeRender[1];
+ if (aWokrSizeY % aLocWorkSize[1] != 0 )
+ aWokrSizeY += aLocWorkSize[1] - aWokrSizeY % aLocWorkSize[1];
+
+ size_t aTotWorkSize[] = { aWorkSizeX, aWokrSizeY };
- size_t aGlbSizeRender[] = { aWorkSizeX, aWokrSizeY };
+ cl_event anEvent = NULL, anEventSmooth = NULL;
+
+ anError = clEnqueueNDRangeKernel (myComputeQueue,
+ myRaytraceRenderKernel, 2, NULL, aTotWorkSize, aLocWorkSize, 0, NULL, &anEvent);
- // Run kernel
- cl_event anEvent (NULL), anEventSmooth (NULL);
- anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceRenderKernel,
- 2, NULL, aGlbSizeRender, aLocSizeRender,
- 0, NULL, &anEvent);
if (anError != CL_SUCCESS)
{
- const TCollection_ExtendedString aMsg = "Error! Failed to execute the ray-tracing kernel!";
+ const TCollection_ExtendedString aMessage = "Error! Failed to execute the ray-tracing kernel!";
+
myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB,
- GL_DEBUG_TYPE_ERROR_ARB,
- 0,
- GL_DEBUG_SEVERITY_HIGH_ARB,
- aMsg);
+ GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+
return Standard_False;
}
+
clWaitForEvents (1, &anEvent);
if (theCView.IsAntialiasingEnabled)
{
- size_t aLocSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4,
- myIsAmdComputePlatform ? 8 : 32 };
+ size_t aLocWorkSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4,
+ myIsAmdComputePlatform ? 8 : 32 };
#ifdef OPENCL_GROUP_SIZE_TEST
- aLocSizeSmooth[0] = aLocX;
- aLocSizeSmooth[1] = aLocY;
+ aLocWorkSizeSmooth[0] = aLocX;
+ aLocWorkSizeSmooth[1] = aLocY;
#endif
aWorkSizeX = theSizeX;
- if (aWorkSizeX % aLocSizeSmooth[0] != 0)
- aWorkSizeX += aLocSizeSmooth[0] - aWorkSizeX % aLocSizeSmooth[0];
+ if (aWorkSizeX % aLocWorkSizeSmooth[0] != 0)
+ aWorkSizeX += aLocWorkSizeSmooth[0] - aWorkSizeX % aLocWorkSizeSmooth[0];
size_t aWokrSizeY = theSizeY;
- if (aWokrSizeY % aLocSizeSmooth[1] != 0 )
- aWokrSizeY += aLocSizeSmooth[1] - aWokrSizeY % aLocSizeSmooth[1];
+ if (aWokrSizeY % aLocWorkSizeSmooth[1] != 0 )
+ aWokrSizeY += aLocWorkSizeSmooth[1] - aWokrSizeY % aLocWorkSizeSmooth[1];
+
+ size_t aTotWorkSizeSmooth [] = { aWorkSizeX, aWokrSizeY };
+
+ anError = clEnqueueNDRangeKernel (myComputeQueue, myRaytraceSmoothKernel,
+ 2, NULL, aTotWorkSizeSmooth, aLocWorkSizeSmooth, 0, NULL, &anEventSmooth);
- size_t aGlbSizeSmooth [] = { aWorkSizeX, aWokrSizeY };
- anError = clEnqueueNDRangeKernel (myRaytraceQueue, myRaytraceSmoothKernel,
- 2, NULL, aGlbSizeSmooth, aLocSizeSmooth,
- 0, NULL, &anEventSmooth);
clWaitForEvents (1, &anEventSmooth);
if (anError != CL_SUCCESS)
{
- const TCollection_ExtendedString aMsg = "Error! Failed to execute the 'smoothing' kernel!";
+ 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,
- aMsg);
+ GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage);
+
return Standard_False;
}
}
- // Get the profiling data
-#if defined (RAY_TRACE_PRINT_INFO) || defined(OPENCL_GROUP_SIZE_TEST)
+#if defined (RAY_TRACE_PRINT_INFO) || defined (OPENCL_GROUP_SIZE_TEST)
+
+ static cl_ulong ttt1 = 10000000000;
+ static cl_ulong ttt2 = 10000000000;
+
+ cl_ulong aBegTime = 0;
+ cl_ulong aEndTime = 0;
- cl_ulong aTimeStart,
- aTimeFinal;
+ clGetEventProfilingInfo (anEvent,
+ CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL);
+ clGetEventProfilingInfo (anEvent,
+ CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL);
- clGetEventProfilingInfo (anEvent, CL_PROFILING_COMMAND_START,
- sizeof(aTimeStart), &aTimeStart, NULL);
- clGetEventProfilingInfo (anEvent, CL_PROFILING_COMMAND_END,
- sizeof(aTimeFinal), &aTimeFinal, NULL);
- std::cout << "\tRender time (ms): " << ( aTimeFinal - aTimeStart ) / 1e6f << std::endl;
+ ttt1 = aEndTime - aBegTime < ttt1 ? aEndTime - aBegTime : ttt1;
+
+ std::cout << "\tRender time (ms): " << ttt1 / 1e6f << std::endl;
if (theCView.IsAntialiasingEnabled)
{
- clGetEventProfilingInfo (anEventSmooth, CL_PROFILING_COMMAND_START,
- sizeof(aTimeStart), &aTimeStart, NULL);
- clGetEventProfilingInfo (anEventSmooth, CL_PROFILING_COMMAND_END,
- sizeof(aTimeFinal), &aTimeFinal, NULL);
- std::cout << "\tSmoothing time (ms): " << ( aTimeFinal - aTimeStart ) / 1e6f << std::endl;
+ clGetEventProfilingInfo (anEventSmooth,
+ CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL);
+ clGetEventProfilingInfo (anEventSmooth,
+ CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL);
+
+ ttt2 = aEndTime - aBegTime < ttt2 ? aEndTime - aBegTime : ttt2;
+
+ std::cout << "\tSmooth time (ms): " << ttt2 / 1e6f << std::endl;
}
+
#endif
if (anEvent != NULL)
if (anEventSmooth != NULL)
clReleaseEvent (anEventSmooth);
+
+#ifdef OPENCL_GROUP_SIZE_TEST
}
+#endif
return Standard_True;
}
{
for (Standard_Integer x = -1; x <= 1; x += 2)
{
- OpenGl_RTVec4f aOrigin (float(x),
- float(y),
- -1.f,
- 1.f);
+ BVH_Vec4f aOrigin (float(x),
+ float(y),
+ -1.f,
+ 1.f);
aOrigin = MatVecMult (theInvModelProj, aOrigin);
aOrigin.x() = aOrigin.x() / aOrigin.w();
aOrigin.z() = aOrigin.z() / aOrigin.w();
aOrigin.w() = 1.f;
- OpenGl_RTVec4f aDirect (float(x),
- float(y),
- 1.f,
- 1.f);
+ BVH_Vec4f aDirect (float(x),
+ float(y),
+ 1.f,
+ 1.f);
aDirect = MatVecMult (theInvModelProj, aDirect);
aDirect.x() = aDirect.x() / aDirect.w();
// Compute ray-traced image using OpenCL kernel
cl_mem anImages[] = { myRaytraceOutputImage, myRaytraceOutputImageAA };
- cl_int anError = clEnqueueAcquireGLObjects (myRaytraceQueue,
+ cl_int anError = clEnqueueAcquireGLObjects (myComputeQueue,
2, anImages,
0, NULL, NULL);
- clFinish (myRaytraceQueue);
+ clFinish (myComputeQueue);
if (myIsRaytraceDataValid)
{
theSizeY);
}
- anError |= clEnqueueReleaseGLObjects (myRaytraceQueue,
+ anError |= clEnqueueReleaseGLObjects (myComputeQueue,
2, anImages,
0, NULL, NULL);
- clFinish (myRaytraceQueue);
+ clFinish (myComputeQueue);
// Draw background
glPushAttrib (GL_ENABLE_BIT |
vdisplay s1 s2
vsetmaterial s1 Silver
vsetmaterial s2 Pewter
+vlight change 0 pos -1 1 1
vfit
# activate ray-tracing
projects / occt.git / commitdiff