// Created on: 2013-10-16
// 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

#define EOL "\n"

extern const char THE_RAY_TRACE_OPENCL_SOURCE[] =

  /////////////////////////////////////////////////////////////////////////////////////////
  // Specific data types
  EOL
  //! Stores ray parameters.
  EOL"  typedef struct __SRay"
  EOL"  {"
  EOL"    float4 Origin;"
  EOL"    float4 Direct;"
  EOL"  }"
  EOL"  SRay;"
  EOL
  //! Stores parameters of intersection point.
  EOL"  typedef struct __SIntersect"
  EOL"  {"
  EOL"    float4 Normal;"
  EOL"    float Time;"
  EOL"    float U;"
  EOL"    float V;"
  EOL"  }"
  EOL"  SIntersect;"
  EOL
  EOL
  /////////////////////////////////////////////////////////////////////////////////////////
  // Some useful vector constants
  EOL
  EOL"  #define ZERO ( float4 )( 0.f, 0.f, 0.f, 0.f )"
  EOL"  #define UNIT ( float4 )( 1.f, 1.f, 1.f, 0.f )"
  EOL
  EOL"  #define AXIS_X ( float4 )( 1.f, 0.f, 0.f, 0.f )"
  EOL"  #define AXIS_Y ( float4 )( 0.f, 1.f, 0.f, 0.f )"
  EOL"  #define AXIS_Z ( float4 )( 0.f, 0.f, 1.f, 0.f )"
  EOL
  /////////////////////////////////////////////////////////////////////////////////////////
  // Support functions
  EOL
  // =======================================================================
  // function : GenerateRay
  // purpose  : Generates primary ray for current work item
  // =======================================================================
  EOL"  void GenerateRay (SRay* theRay,"
  EOL"                    const float theX,"
  EOL"                    const float theY,"
  EOL"                    const int theSizeX,"
  EOL"                    const int theSizeY,"
  EOL"                    const float16 theOrigins,"
  EOL"                    const float16 theDirects)"
  EOL"  {"
  EOL"    float2 aPixel = (float2) (theX / (float)theSizeX,"
  EOL"                              theY / (float)theSizeY);"
  EOL
  EOL"    float4 aP0 = mix (theOrigins.lo.lo, theOrigins.lo.hi, aPixel.x);"
  EOL"    float4 aP1 = mix (theOrigins.hi.lo, theOrigins.hi.hi, aPixel.x);"
  EOL
  EOL"    theRay->Origin = mix (aP0, aP1, aPixel.y);"
  EOL
  EOL"    aP0 = mix (theDirects.lo.lo, theDirects.lo.hi, aPixel.x);"
  EOL"    aP1 = mix (theDirects.hi.lo, theDirects.hi.hi, aPixel.x);"
  EOL
  EOL"    theRay->Direct = mix (aP0, aP1, aPixel.y);"
  EOL"  }"
  EOL
  EOL
  /////////////////////////////////////////////////////////////////////////////////////////
  // Functions for compute ray-object intersection
  EOL
  EOL"  #define _OOEPS_ exp2 (-80.0f)"
  EOL
  // =======================================================================
  // function : IntersectSphere
  // purpose  : Computes ray-sphere intersection
  // =======================================================================
  EOL"  float IntersectSphere (const SRay* theRay, float theRadius)"
  EOL"  {"
  EOL"    float aDdotD = dot (theRay->Direct.xyz, theRay->Direct.xyz);"
  EOL"    float aDdotO = dot (theRay->Direct.xyz, theRay->Origin.xyz);"
  EOL"    float aOdotO = dot (theRay->Origin.xyz, theRay->Origin.xyz);"
  EOL
  EOL"    float aD = aDdotO * aDdotO - aDdotD * (aOdotO - theRadius * theRadius);"
  EOL
  EOL"    if (aD > 0.f)"
  EOL"    {"
  EOL"      float aTime = (-aDdotO + native_sqrt (aD)) * (1.f / aDdotD);"
  EOL
  EOL"      return aTime > 0.f ? aTime : MAXFLOAT;"
  EOL"    }"
  EOL
  EOL"    return MAXFLOAT;"
  EOL"  }"
  EOL
  // =======================================================================
  // function : IntersectBox
  // purpose  : Computes ray-box intersection (slab test)
  // =======================================================================
  EOL"  float IntersectBox (const SRay* theRay,"
  EOL"                      float4 theMinPoint,"
  EOL"                      float4 theMaxPoint)"
  EOL"  {"
  EOL"    const float4 aInvDirect = (float4)("
  EOL"                    1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
  EOL"                           theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
  EOL"                    1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
  EOL"                           theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
  EOL"                    1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
  EOL"                           theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
  EOL"                    0.f);"
  EOL
  EOL"    const float4 aTime0 = (theMinPoint - theRay->Origin) * aInvDirect;"
  EOL"    const float4 aTime1 = (theMaxPoint - theRay->Origin) * aInvDirect;"
  EOL
  EOL"    const float4 aTimeMax = max (aTime0, aTime1);"
  EOL"    const float4 aTimeMin = min (aTime0, aTime1);"
  EOL
  EOL"    const float theTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
  EOL"    const float theTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
  EOL
  EOL"    return (theTimeStart <= theTimeFinal) && (theTimeFinal >= 0.f) ? theTimeStart : MAXFLOAT;"
  EOL"  }"
  EOL
  // =======================================================================
  // function : IntersectNodes
  // purpose  : Computes intersection of ray with two child nodes (boxes)
  // =======================================================================
  EOL"  void IntersectNodes (const SRay* theRay,"
  EOL"                       float4 theMinPoint0,"
  EOL"                       float4 theMaxPoint0,"
  EOL"                       float4 theMinPoint1,"
  EOL"                       float4 theMaxPoint1,"
  EOL"                       float* theTimeStart0,"
  EOL"                       float* theTimeStart1,"
  EOL"                       float theMaxTime)"
  EOL"  {"
  EOL"    const float4 aInvDirect = (float4)("
  EOL"                    1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
  EOL"                           theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
  EOL"                    1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
  EOL"                           theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
  EOL"                    1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
  EOL"                           theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
  EOL"                    0.f);"
  EOL
  EOL"    float4 aTime0 = (theMinPoint0 - theRay->Origin) * aInvDirect;"
  EOL"    float4 aTime1 = (theMaxPoint0 - theRay->Origin) * aInvDirect;"
  EOL
  EOL"    float4 aTimeMax = max (aTime0, aTime1);"
  EOL"    float4 aTimeMin = min (aTime0, aTime1);"
  EOL
  EOL"    aTime0 = (theMinPoint1 - theRay->Origin) * aInvDirect;"
  EOL"    aTime1 = (theMaxPoint1 - theRay->Origin) * aInvDirect;"
  EOL
  EOL"    float aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
  EOL"    float aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
  EOL
  EOL"    aTimeMax = max (aTime0, aTime1);"
  EOL"    aTimeMin = min (aTime0, aTime1);"
  EOL
  EOL"    *theTimeStart0 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
  EOL"                   ? aTimeStart : -MAXFLOAT;"
  EOL
  EOL"    aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
  EOL"    aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
  EOL
  EOL"    *theTimeStart1 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
  EOL"                   ? aTimeStart : -MAXFLOAT;"
  EOL"  }"
  EOL
  // =======================================================================
  // function : IntersectTriangle
  // purpose  : Computes ray-triangle intersection (branchless version)
  // =======================================================================
  EOL"   float IntersectTriangle (const SRay* theRay,"
  EOL"                            const float4 thePoint0,"
  EOL"                            const float4 thePoint1,"
  EOL"                            const float4 thePoint2,"
  EOL"                            float4* theNormal,"
  EOL"                            float* theU,"
  EOL"                            float* theV)"
  EOL"  {"
  EOL"    const float4 aEdge0 = thePoint1 - thePoint0;"
  EOL"    const float4 aEdge1 = thePoint0 - thePoint2;"
  EOL
  EOL"    *theNormal = cross (aEdge1, aEdge0);"
  EOL
  EOL"    const float4 aEdge2 = (1.f / dot (*theNormal, theRay->Direct)) * (thePoint0 - theRay->Origin);"
  EOL
  EOL"    const float aTime = dot (*theNormal, aEdge2);"
  EOL
  EOL"    const float4 theVec = cross (theRay->Direct, aEdge2);"
  EOL
  EOL"    *theU = dot (theVec, aEdge1);"
  EOL"    *theV = dot (theVec, aEdge0);"
  EOL
  EOL"    return (aTime >= 0.f) & (*theU >= 0.f) & (*theV >= 0.f) & (*theU + *theV <= 1.f) ? aTime : MAXFLOAT;"
  EOL"  }"
  EOL
  /////////////////////////////////////////////////////////////////////////////////////////
  // Support shading functions
  EOL
  EOL"  const sampler_t EnvironmentSampler ="
  EOL"            CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_REPEAT | CLK_FILTER_LINEAR;"
  EOL
  // =======================================================================
  // function : SmoothNormal
  // purpose  : Interpolates normal across the triangle
  // =======================================================================
  EOL"  float4 SmoothNormal (__global float4* theNormals,"
  EOL"                       const SIntersect* theHit,"
  EOL"                       const int4 theIndices)"
  EOL"  {"
  EOL"    float4 aNormal0 = theNormals[theIndices.x],"
  EOL"           aNormal1 = theNormals[theIndices.y],"
  EOL"           aNormal2 = theNormals[theIndices.z];"
  EOL
  EOL"    return fast_normalize (aNormal1 * theHit->U +"
  EOL"                           aNormal2 * theHit->V +"
  EOL"                           aNormal0 * (1.f - theHit->U - theHit->V));"
  EOL"  }"
  EOL
  // =======================================================================
  // function : Shade
  // purpose  : Computes Phong-based illumination
  // =======================================================================
  EOL"  float4 Shade (const float4 theMatDiff,"
  EOL"                const float4 theMatSpec,"
  EOL"                const float4 theLight,"
  EOL"                const float4 theView,"
  EOL"                const float4 theNormal,"
  EOL"                const float4 theIntens,"
  EOL"                const float theTranspr)"
  EOL"  {"
  EOL"    float aLambert = dot (theNormal, theLight);"
  EOL
  EOL"    aLambert = theTranspr > 0.f ? fabs (aLambert) : aLambert;"
  EOL
  EOL"    if (aLambert > 0.f)"
  EOL"    {"
  EOL"      const float4 aReflect = 2.f * dot (theLight, theNormal) * theNormal - theLight;"
  EOL
  EOL"      const float aSpecular = pow (max (dot (aReflect.xyz, theView.xyz), 0.f), theMatSpec.w);"
  EOL
  EOL"      return theIntens * (theMatDiff * aLambert + theMatSpec * aSpecular);"
  EOL"    }"
  EOL
  EOL"    return ZERO;"
  EOL"  }"
  EOL
  // =======================================================================
  // function : Latlong
  // purpose  : Converts world direction to environment texture coordinates
  // =======================================================================
  EOL"  float2 Latlong (const float4 thePoint, const float theRadius)"
  EOL"  {"
  EOL"    float aPsi = acospi (-thePoint.y / theRadius);"
  EOL"    float aPhi = atan2pi (thePoint.z, thePoint.x);"
  EOL
  EOL"    aPhi = (aPhi < 0.f) ? aPhi + 2.f : aPhi;"
  EOL
  EOL"    return (float2) (aPhi * 0.5f, aPsi);"
  EOL"  }"
  EOL
  /////////////////////////////////////////////////////////////////////////////////////////
  // Core ray tracing function
  EOL
  // =======================================================================
  // function : push
  // purpose  : Pushes BVH node index to local stack
  // =======================================================================
  EOL"  void push (uint* theStack, char* thePos, const uint theValue)"
  EOL"  {"
  EOL"    (*thePos)++;"
  EOL"    theStack[*thePos] = theValue;"
  EOL"  }"
  EOL
  // =======================================================================
  // function : pop
  // purpose  : Pops BVH node index from local stack
  // =======================================================================
  EOL"  void pop (uint* theStack, char* thePos, uint* theValue)"
  EOL"  {"
  EOL"    *theValue = theStack[*thePos];"
  EOL"    (*thePos)--;"
  EOL"  }"
  EOL
  // =======================================================================
  // function : ObjectNearestHit
  // purpose  : Finds intersection with nearest object triangle
  // =======================================================================
  EOL"  int4 ObjectNearestHit (const SRay* theRay,"
  EOL"                         SIntersect* theIntersect,"
  EOL"                         __global int4* theObjectNodeInfoBuffer,"
  EOL"                         __global float4* theObjectMinPointBuffer,"
  EOL"                         __global float4* theObjectMaxPointBuffer,"
  EOL"                         __global int4* theGeometryTriangBuffer,"
  EOL"                         __global float4* theGeometryVertexBuffer)"
  EOL"  {"
  EOL"    uint aStack [32];"
  EOL
  EOL"    char aHead = -1;" // stack pointer
  EOL"    uint aNode =  0;" // node to visit
  EOL
  EOL"    const float4 aInvDirect = (float4) ("
  EOL"      1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
  EOL"             theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
  EOL"      1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
  EOL"             theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
  EOL"      1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
  EOL"             theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
  EOL"      0.f);"
  EOL
  EOL"    int4 aTriangleIndex = (int4) (-1);"
  EOL
  EOL"    float aTimeExit;"
  EOL"    float aTimeMin1;"
  EOL"    float aTimeMin2;"
  EOL
  EOL"    while (true)"
  EOL"    {"
  EOL"      const int3 aData = theObjectNodeInfoBuffer[aNode].xyz;"
  EOL
  EOL"      if (aData.x == 0)" // if inner node
  EOL"      {"
  EOL"        float4 aNodeMin = theObjectMinPointBuffer[aData.y];"
  EOL"        float4 aNodeMax = theObjectMaxPointBuffer[aData.y];"
  EOL
  EOL"        float4 aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
  EOL"        float4 aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
  EOL
  EOL"        float4 aTimeMax = max (aTime0, aTime1);"
  EOL"        float4 aTimeMin = min (aTime0, aTime1);"
  EOL
  EOL"        aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
  EOL"        aTimeMin1 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
  EOL
  EOL"        const bool aHitLft = (aTimeMin1 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin1 <= theIntersect->Time);"
  EOL
  EOL"        aNodeMin = theObjectMinPointBuffer[aData.z];"
  EOL"        aNodeMax = theObjectMaxPointBuffer[aData.z];"
  EOL
  EOL"        aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
  EOL"        aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
  EOL
  EOL"        aTimeMax = max (aTime0, aTime1);"
  EOL"        aTimeMin = min (aTime0, aTime1);"
  EOL
  EOL"        aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
  EOL"        aTimeMin2 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
  EOL
  EOL"        const bool aHitRgh = (aTimeMin2 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin2 <= theIntersect->Time);"
  EOL
  EOL"        if (aHitLft & aHitRgh)"
  EOL"        {"
  EOL"          aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
  EOL
  EOL"          push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
  EOL"        }"
  EOL"        else"
  EOL"        {"
  EOL"          if (aHitLft | aHitRgh)"
  EOL"          {"
  EOL"            aNode = aHitLft ? aData.y : aData.z;"
  EOL"          }"
  EOL"          else"
  EOL"          {"
  EOL"            if (aHead < 0)"
  EOL"              return aTriangleIndex;"
  EOL
  EOL"            pop (aStack, &aHead, &aNode);"
  EOL"          }"
  EOL"        }"
  EOL"      }"
  EOL"      else " // if leaf node
  EOL"      {"
  EOL"        for (int aTriIdx = aData.y; aTriIdx <= aData.z; ++aTriIdx)"
  EOL"        {"
  EOL"          const int4 aTestTriangle = theGeometryTriangBuffer[aTriIdx];"
  EOL
  EOL"          const float4 aPoint0 = theGeometryVertexBuffer[aTestTriangle.x];"
  EOL"          const float4 aPoint1 = theGeometryVertexBuffer[aTestTriangle.y];"
  EOL"          const float4 aPoint2 = theGeometryVertexBuffer[aTestTriangle.z];"
  EOL
  EOL"          float4 aNormal; float aU, aV;"
  EOL
  EOL"          float aTime = IntersectTriangle (theRay,"
  EOL"                                           aPoint0,"
  EOL"                                           aPoint1,"
  EOL"                                           aPoint2,"
  EOL"                                           &aNormal,"
  EOL"                                           &aU,"
  EOL"                                           &aV);"
  EOL
  EOL"          if (aTime < theIntersect->Time)"
  EOL"          {"
  EOL"            aTriangleIndex = aTestTriangle;"
  EOL"            theIntersect->Normal = aNormal;"
  EOL"            theIntersect->Time = aTime;"
  EOL"            theIntersect->U = aU;"
  EOL"            theIntersect->V = aV;"
  EOL"          }"
  EOL"        }"
  EOL
  EOL"        if (aHead < 0)"
  EOL"          return aTriangleIndex;"
  EOL
  EOL"        pop (aStack, &aHead, &aNode);"
  EOL"      }"
  EOL"    }"
  EOL
  EOL"    return aTriangleIndex;"
  EOL"  }"
  EOL
  // =======================================================================
  // function : ObjectAnyHit
  // purpose  : Finds intersection with any object triangle
  // =======================================================================
  EOL"  float ObjectAnyHit (const SRay* theRay,"
  EOL"                      __global int4* theObjectNodeInfoBuffer,"
  EOL"                      __global float4* theObjectMinPointBuffer,"
  EOL"                      __global float4* theObjectMaxPointBuffer,"
  EOL"                      __global int4* theGeometryTriangBuffer,"
  EOL"                      __global float4* theGeometryVertexBuffer,"
  EOL"                      const float theDistance)"
  EOL"  {"
  EOL"    uint aStack [32];"
  EOL
  EOL"    char aHead = -1;" // stack pointer
  EOL"    uint aNode =  0;" // node to visit
  EOL
  EOL"    const float4 aInvDirect = (float4) ("
  EOL"      1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
  EOL"             theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
  EOL"      1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
  EOL"             theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
  EOL"      1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
  EOL"             theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
  EOL"      0.f);"
  EOL
  EOL"    float aTimeExit;"
  EOL"    float aTimeMin1;"
  EOL"    float aTimeMin2;"
  EOL
  EOL"    while (true)"
  EOL"    {"
  EOL"      const int3 aData = theObjectNodeInfoBuffer[aNode].xyz;"
  EOL
  EOL"      if (aData.x == 0)" // if inner node
  EOL"      {"
  EOL"        float4 aNodeMin = theObjectMinPointBuffer[aData.y];"
  EOL"        float4 aNodeMax = theObjectMaxPointBuffer[aData.y];"
  EOL
  EOL"        float4 aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
  EOL"        float4 aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
  EOL
  EOL"        float4 aTimeMax = max (aTime0, aTime1);"
  EOL"        float4 aTimeMin = min (aTime0, aTime1);"
  EOL
  EOL"        aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
  EOL"        aTimeMin1 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
  EOL
  EOL"        const bool aHitLft = (aTimeMin1 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin1 <= theDistance);"
  EOL
  EOL"        aNodeMin = theObjectMinPointBuffer[aData.z];"
  EOL"        aNodeMax = theObjectMaxPointBuffer[aData.z];"
  EOL
  EOL"        aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
  EOL"        aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
  EOL
  EOL"        aTimeMax = max (aTime0, aTime1);"
  EOL"        aTimeMin = min (aTime0, aTime1);"
  EOL
  EOL"        aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
  EOL"        aTimeMin2 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
  EOL
  EOL"        const bool aHitRgh = (aTimeMin2 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin2 <= theDistance);"
  EOL
  EOL"        if (aHitLft & aHitRgh)"
  EOL"        {"
  EOL"          aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
  EOL
  EOL"          push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
  EOL"        }"
  EOL"        else"
  EOL"        {"
  EOL"          if (aHitLft | aHitRgh)"
  EOL"          {"
  EOL"            aNode = aHitLft ? aData.y : aData.z;"
  EOL"          }"
  EOL"          else"
  EOL"          {"
  EOL"            if (aHead < 0)"
  EOL"              return 1.f;"
  EOL
  EOL"            pop (aStack, &aHead, &aNode);"
  EOL"          }"
  EOL"        }"
  EOL"      }"
  EOL"      else " // if leaf node
  EOL"      {"
  EOL"        for (int aTriIdx = aData.y; aTriIdx <= aData.z; ++aTriIdx)"
  EOL"        {"
  EOL"          const int4 aTestTriangle = theGeometryTriangBuffer[aTriIdx];"
  EOL
  EOL"          const float4 aPoint0 = theGeometryVertexBuffer[aTestTriangle.x];"
  EOL"          const float4 aPoint1 = theGeometryVertexBuffer[aTestTriangle.y];"
  EOL"          const float4 aPoint2 = theGeometryVertexBuffer[aTestTriangle.z];"
  EOL
  EOL"          float4 aNormal; float aU, aV;"
  EOL
  EOL"          float aTime = IntersectTriangle (theRay,"
  EOL"                                           aPoint0,"
  EOL"                                           aPoint1,"
  EOL"                                           aPoint2,"
  EOL"                                           &aNormal,"
  EOL"                                           &aU,"
  EOL"                                           &aV);"
  EOL
  EOL"          if (aTime < theDistance)"
  EOL"          {"
  EOL"            return 0.f;"
  EOL"          }"
  EOL"        }"
  EOL
  EOL"        if (aHead < 0)"
  EOL"          return 1.f;"
  EOL
  EOL"        pop (aStack, &aHead, &aNode);"
  EOL"      }"
  EOL"    }"
  EOL
  EOL"    return 1.f;"
  EOL"  }"
  EOL
  // =======================================================================
  // function : NearestHit
  // purpose  : Finds intersection with nearest scene triangle
  // =======================================================================
  EOL"  int4 NearestHit (const SRay* theRay,"
  EOL"                   SIntersect* theIntersect,"
  EOL"                   __global int4* theSceneNodeInfoBuffer,"
  EOL"                   __global float4* theSceneMinPointBuffer,"
  EOL"                   __global float4* theSceneMaxPointBuffer,"
  EOL"                   __global int4* theObjectNodeInfoBuffer,"
  EOL"                   __global float4* theObjectMinPointBuffer,"
  EOL"                   __global float4* theObjectMaxPointBuffer,"
  EOL"                   __global int4* theGeometryTriangBuffer,"
  EOL"                   __global float4* theGeometryVertexBuffer)"
  EOL"  {"
  EOL"    theIntersect->Time = MAXFLOAT;"
  EOL
  EOL"    uint aStack [16];"
  EOL
  EOL"    char aHead = -1;" // stack pointer
  EOL"    uint aNode =  0;" // node to visit
  EOL
  EOL"    int4 aNearestTriangle = (int4) (-1);"
  EOL
  EOL"    while (true)"
  EOL"    {"
  EOL"      const int4 aData = theSceneNodeInfoBuffer[aNode];"
  EOL
  EOL"      if (aData.x != 0)" // if leaf node
  EOL"      {"
  EOL"        const float4 aNodeMin = theSceneMinPointBuffer[aNode];"
  EOL"        const float4 aNodeMax = theSceneMaxPointBuffer[aNode];"
  EOL
  EOL"        if (IntersectBox (theRay, aNodeMin, aNodeMax) <= theIntersect->Time)"
  EOL"        {"
  EOL"          int4 anIndex = ObjectNearestHit (theRay,"
  EOL"                                           theIntersect,"
  EOL"                                           theObjectNodeInfoBuffer + aData.y,"
  EOL"                                           theObjectMinPointBuffer + aData.y,"
  EOL"                                           theObjectMaxPointBuffer + aData.y,"
  EOL"                                           theGeometryTriangBuffer + aData.w,"
  EOL"                                           theGeometryVertexBuffer + aData.z);"
  EOL
  EOL"          if (anIndex.x != -1)"
  EOL"            aNearestTriangle = (int4) (anIndex.x + aData.z,"
  EOL"                                       anIndex.y + aData.z,"
  EOL"                                       anIndex.z + aData.z,"
  EOL"                                       anIndex.w);"
  EOL"        }"
  EOL
  EOL"        if (aHead < 0)"
  EOL"          return aNearestTriangle;"
  EOL
  EOL"        pop (aStack, &aHead, &aNode);"
  EOL"      }"
  EOL"      else " // if inner node
  EOL"      {"
  EOL"        float4 aNodeMinLft = theSceneMinPointBuffer[aData.y];"
  EOL"        float4 aNodeMinRgh = theSceneMinPointBuffer[aData.z];"
  EOL"        float4 aNodeMaxLft = theSceneMaxPointBuffer[aData.y];"
  EOL"        float4 aNodeMaxRgh = theSceneMaxPointBuffer[aData.z];"
  EOL
  EOL"        float aTimeMin1;"
  EOL"        float aTimeMin2;"
  EOL
  EOL"        IntersectNodes (theRay,"
  EOL"                        aNodeMinLft,"
  EOL"                        aNodeMaxLft,"
  EOL"                        aNodeMinRgh,"
  EOL"                        aNodeMaxRgh,"
  EOL"                        &aTimeMin1,"
  EOL"                        &aTimeMin2,"
  EOL"                        theIntersect->Time);"
  EOL
  EOL"        const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
  EOL"        const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
  EOL
  EOL"        if (aHitLft & aHitRgh)"
  EOL"        {"
  EOL"          aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
  EOL
  EOL"          push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
  EOL"        }"
  EOL"        else"
  EOL"        {"
  EOL"          if (aHitLft | aHitRgh)"
  EOL"          {"
  EOL"            aNode = aHitLft ? aData.y : aData.z;"
  EOL"          }"
  EOL"          else"
  EOL"          {"
  EOL"            if (aHead < 0)"
  EOL"              return aNearestTriangle;"
  EOL
  EOL"            pop (aStack, &aHead, &aNode);"
  EOL"          }"
  EOL"        }"
  EOL"      }"
  EOL"    }"
  EOL
  EOL"    return aNearestTriangle;"
  EOL"  }"
  EOL
  // =======================================================================
  // function : AnyHit
  // purpose  : Finds intersection with any scene triangle
  // =======================================================================
  EOL"  float AnyHit (const SRay* theRay,"
  EOL"                __global int4* theSceneNodeInfoBuffer,"
  EOL"                __global float4* theSceneMinPointBuffer,"
  EOL"                __global float4* theSceneMaxPointBuffer,"
  EOL"                __global int4* theObjectNodeInfoBuffer,"
  EOL"                __global float4* theObjectMinPointBuffer,"
  EOL"                __global float4* theObjectMaxPointBuffer,"
  EOL"                __global int4* theGeometryTriangBuffer,"
  EOL"                __global float4* theGeometryVertexBuffer,"
  EOL"                const float theDistance)"
  EOL"  {"
  EOL"    uint aStack [16];"
  EOL
  EOL"    char aHead = -1;" // stack pointer
  EOL"    uint aNode =  0;" // node to visit
  EOL
  EOL"    while (true)"
  EOL"    {"
  EOL"      const int4 aData = theSceneNodeInfoBuffer[aNode];"
  EOL
  EOL"      if (aData.x != 0)" // if leaf node
  EOL"      {"
  EOL"        const float4 aNodeMin = theSceneMinPointBuffer[aNode];"
  EOL"        const float4 aNodeMax = theSceneMaxPointBuffer[aNode];"
  EOL
  EOL"        if (IntersectBox (theRay, aNodeMin, aNodeMax) <= theDistance)"
  EOL"        {"
  EOL"          if (0.f == ObjectAnyHit (theRay,"
  EOL"                                   theObjectNodeInfoBuffer + aData.y,"
  EOL"                                   theObjectMinPointBuffer + aData.y,"
  EOL"                                   theObjectMaxPointBuffer + aData.y,"
  EOL"                                   theGeometryTriangBuffer + aData.w,"
  EOL"                                   theGeometryVertexBuffer + aData.z,"
  EOL"                                   theDistance))"
  EOL"          {"
  EOL"            return 0.f;"
  EOL"          }"
  EOL"        }"
  EOL
  EOL"        if (aHead < 0)"
  EOL"          return 1.f;"
  EOL
  EOL"        pop (aStack, &aHead, &aNode);"
  EOL"      }"
  EOL"      else" // if inner node
  EOL"      {"
  EOL"        float4 aNodeMinLft = theSceneMinPointBuffer[aData.y];"
  EOL"        float4 aNodeMinRgh = theSceneMinPointBuffer[aData.z];"
  EOL"        float4 aNodeMaxLft = theSceneMaxPointBuffer[aData.y];"
  EOL"        float4 aNodeMaxRgh = theSceneMaxPointBuffer[aData.z];"
  EOL
  EOL"        float aTimeMin1;"
  EOL"        float aTimeMin2;"
  EOL
  EOL"        IntersectNodes (theRay,"
  EOL"                        aNodeMinLft,"
  EOL"                        aNodeMaxLft,"
  EOL"                        aNodeMinRgh,"
  EOL"                        aNodeMaxRgh,"
  EOL"                        &aTimeMin1,"
  EOL"                        &aTimeMin2,"
  EOL"                        theDistance);"
  EOL
  EOL"        const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
  EOL"        const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
  EOL
  EOL"        if (aHitLft & aHitRgh)"
  EOL"        {"
  EOL"          aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
  EOL
  EOL"          push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
  EOL"        }"
  EOL"        else"
  EOL"        {"
  EOL"          if (aHitLft | aHitRgh)"
  EOL"          {"
  EOL"            aNode = aHitLft ? aData.y : aData.z;"
  EOL"          }"
  EOL"          else"
  EOL"          {"
  EOL"            if (aHead < 0)"
  EOL"              return 1.f;"
  EOL
  EOL"            pop (aStack, &aHead, &aNode);"
  EOL"          }"
  EOL"        }"
  EOL"      }"
  EOL"    }"
  EOL"  }"
  EOL
  EOL"  #define _MAX_DEPTH_ 5"
  EOL
  EOL"  #define THRESHOLD (float4) (0.1f, 0.1f, 0.1f, 1.f)"
  EOL
  EOL"  #define LIGHT_POS(Buffer, LightID) Buffer[2 * LightID + 1]"
  EOL"  #define LIGHT_RAD(Buffer, LightID) Buffer[2 * LightID + 0]"
  EOL
  EOL"  #define MATERIAL_AMBN(Buffer, TriangleID) Buffer[7 * TriangleID.w + 0]"
  EOL"  #define MATERIAL_DIFF(Buffer, TriangleID) Buffer[7 * TriangleID.w + 1]"
  EOL"  #define MATERIAL_SPEC(Buffer, TriangleID) Buffer[7 * TriangleID.w + 2]"
  EOL"  #define MATERIAL_EMIS(Buffer, TriangleID) Buffer[7 * TriangleID.w + 3]"
  EOL"  #define MATERIAL_REFL(Buffer, TriangleID) Buffer[7 * TriangleID.w + 4]"
  EOL"  #define MATERIAL_REFR(Buffer, TriangleID) Buffer[7 * TriangleID.w + 5]"
  EOL"  #define MATERIAL_TRAN(Buffer, TriangleID) Buffer[7 * TriangleID.w + 6]"
  EOL
  // =======================================================================
  // function : Radiance
  // purpose  : Computes color of specified ray
  // =======================================================================
  EOL"  float4 Radiance (SRay* theRay,"
  EOL"                   __read_only image2d_t theEnvMap,"
  EOL"                   __global int4* theSceneNodeInfoBuffer,"
  EOL"                   __global float4* theSceneMinPointBuffer,"
  EOL"                   __global float4* theSceneMaxPointBuffer,"
  EOL"                   __global int4* theObjectNodeInfoBuffer,"
  EOL"                   __global float4* theObjectMinPointBuffer,"
  EOL"                   __global float4* theObjectMaxPointBuffer,"
  EOL"                   __global int4* theGeometryTriangBuffer,"
  EOL"                   __global float4* theGeometryVertexBuffer,"
  EOL"                   __global float4* theGeometryNormalBuffer,"
  EOL"                   __global float4* theLightSourceBuffer,"
  EOL"                   __global float4* theMaterialBuffer,"
  EOL"                   const float4 theGlobalAmbient,"
  EOL"                   const int theLightBufferSize,"
  EOL"                   const int theShadowsEnabled,"
  EOL"                   const int theReflectEnabled,"
  EOL"                   const float theSceneEpsilon,"
  EOL"                   const float theSceneRadius)"
  EOL"  {"
  EOL"    float4 aResult = (float4) (0.f, 0.f, 0.f, 0.f);"
  EOL"    float4 aWeight = (float4) (1.f, 1.f, 1.f, 1.f);"
  EOL
  EOL"    SIntersect aHit;"
  EOL
  EOL"    for (int aDepth = 0; aDepth < _MAX_DEPTH_; ++aDepth)"
  EOL"    {"
  EOL"      int4 aTriangle = NearestHit (theRay,"
  EOL"                                   &aHit,"
  EOL"                                   theSceneNodeInfoBuffer,"
  EOL"                                   theSceneMinPointBuffer,"
  EOL"                                   theSceneMaxPointBuffer,"
  EOL"                                   theObjectNodeInfoBuffer,"
  EOL"                                   theObjectMinPointBuffer,"
  EOL"                                   theObjectMaxPointBuffer,"
  EOL"                                   theGeometryTriangBuffer,"
  EOL"                                   theGeometryVertexBuffer);"
  EOL
  EOL"      if (aTriangle.x < 0.f)"
  EOL"      {"
  EOL"        if (aWeight.w != 0.f)"
  EOL"          break;"
  EOL
  EOL"        float aTime = IntersectSphere (theRay, theSceneRadius);"
  EOL
  EOL"        if (aTime != MAXFLOAT)"
  EOL"        {"
  EOL"          aResult += aWeight * read_imagef (theEnvMap, EnvironmentSampler,"
  EOL"            Latlong (theRay->Origin + theRay->Direct * aTime, theSceneRadius));"
  EOL"        }"
  EOL
  EOL"        return (float4) (aResult.x,"
  EOL"                         aResult.y,"
  EOL"                         aResult.z,"
  EOL"                         aWeight.w);"
  EOL"      }"
  EOL
  EOL       // Compute geometric normal
  EOL"      float4 aGeomNormal = aHit.Normal; aGeomNormal = fast_normalize (aGeomNormal);"
  EOL
  EOL       // Compute interpolated normal
  EOL"      float4 aNormal = SmoothNormal (theGeometryNormalBuffer, &aHit, aTriangle);"
  EOL
  EOL       // Compute intersection point
  EOL"      float4 aPoint = theRay->Direct * aHit.Time + theRay->Origin;"
  EOL
  EOL"      float4 aMaterAmb = MATERIAL_AMBN (theMaterialBuffer, aTriangle);"
  EOL"      float4 aMaterTrn = MATERIAL_TRAN (theMaterialBuffer, aTriangle);"
  EOL
  EOL"      aResult += aWeight * theGlobalAmbient * aMaterAmb *"
  EOL"            (aMaterTrn.x * max (fabs (dot (theRay->Direct, aNormal)), 0.5f));"
  EOL
  EOL"      for (int nLight = 0; nLight < theLightBufferSize; ++nLight)"
  EOL"      {"
  EOL"        float4 aLightPosition = LIGHT_POS (theLightSourceBuffer, nLight);"
  EOL
  EOL"        SRay aShadow;"
  EOL"        aShadow.Direct = aLightPosition;"
  EOL
  EOL"        float aLightDistance = MAXFLOAT;"
  EOL"        if (aLightPosition.w != 0.f)"
  EOL"        {"
  EOL"          aLightDistance = length (aLightPosition - aPoint);"
  EOL"          aShadow.Direct = (aLightPosition - aPoint) * (1.f / aLightDistance);"
  EOL"        }"
  EOL
  EOL"        aShadow.Origin = aPoint + aShadow.Direct * theSceneEpsilon +"
  EOL"                    aGeomNormal * copysign (theSceneEpsilon, dot (aGeomNormal, aShadow.Direct));"
  EOL
  EOL"        float aVisibility = 1.f;"
  EOL
  EOL"        if (theShadowsEnabled)"
  EOL"        {"
  EOL"          aVisibility = AnyHit (&aShadow,"
  EOL"                                theSceneNodeInfoBuffer,"
  EOL"                                theSceneMinPointBuffer,"
  EOL"                                theSceneMaxPointBuffer,"
  EOL"                                theObjectNodeInfoBuffer,"
  EOL"                                theObjectMinPointBuffer,"
  EOL"                                theObjectMaxPointBuffer,"
  EOL"                                theGeometryTriangBuffer,"
  EOL"                                theGeometryVertexBuffer,"
  EOL"                                aLightDistance);"
  EOL"        }"
  EOL
  EOL"        if (aVisibility > 0.f)"
  EOL"        {"
  EOL"          aResult += aMaterTrn.x * aWeight * Shade (MATERIAL_DIFF (theMaterialBuffer, aTriangle),"
  EOL"                                                    MATERIAL_SPEC (theMaterialBuffer, aTriangle),"
  EOL"                                                    aShadow.Direct,"
  EOL"                                                    -theRay->Direct,"
  EOL"                                                    aNormal,"
  EOL"                                                    LIGHT_RAD (theLightSourceBuffer, nLight),"
  EOL"                                                    aMaterTrn.y);"
  EOL"        }"
  EOL"      }"
  EOL
  EOL"      if (aMaterTrn.y > 0.f)"
  EOL"      {"
  EOL"        aWeight *= aMaterTrn.y;"
  EOL"      }"
  EOL"      else"
  EOL"      {"
  EOL"        aWeight *= theReflectEnabled ? MATERIAL_REFL (theMaterialBuffer, aTriangle) : ZERO;"
  EOL
  EOL"        float4 aDirect = theRay->Direct - 2.f * dot (theRay->Direct, aNormal) * aNormal;"
  EOL
  EOL"        float aDdotN = dot (aDirect, aGeomNormal);"
  EOL"        if (aDdotN < 0.f)"
  EOL"          theRay->Direct -= 2.f * dot (theRay->Direct, aGeomNormal) * aGeomNormal;"
  EOL"        else"
  EOL"          theRay->Direct = aDirect;"
  EOL"      }"
  EOL
  EOL"      if (all (islessequal (aWeight, THRESHOLD)))"
  EOL"      {"
  EOL"        return (float4) (aResult.x,"
  EOL"                         aResult.y,"
  EOL"                         aResult.z,"
  EOL"                         aWeight.w);"
  EOL"      }"
  EOL
  EOL"      theRay->Origin = theRay->Direct * theSceneEpsilon + aPoint;"
  EOL"    }"
  EOL
  EOL"    return (float4) (aResult.x,"
  EOL"                     aResult.y,"
  EOL"                     aResult.z,"
  EOL"                     aWeight.w);"
  EOL"  }"
  EOL
  ///////////////////////////////////////////////////////////////////////////////
  // Ray tracing kernel functions
  EOL
  // =======================================================================
  // function : RaytraceRender
  // purpose  : Computes pixel color using ray-tracing
  // =======================================================================
  EOL"  __kernel void RaytraceRender (const int theSizeX,"
  EOL"                                const int theSizeY,"
  EOL"                                const float16 theOrigins,"
  EOL"                                const float16 theDirects,"
  EOL"                                __read_only image2d_t theEnvMap,"
  EOL"                                __write_only image2d_t theOutput,"
  EOL"                                __global int4* theSceneNodeInfoBuffer,"
  EOL"                                __global float4* theSceneMinPointBuffer,"
  EOL"                                __global float4* theSceneMaxPointBuffer,"
  EOL"                                __global int4* theObjectNodeInfoBuffer,"
  EOL"                                __global float4* theObjectMinPointBuffer,"
  EOL"                                __global float4* theObjectMaxPointBuffer,"
  EOL"                                __global int4* theGeometryTriangBuffer,"
  EOL"                                __global float4* theGeometryVertexBuffer,"
  EOL"                                __global float4* theGeometryNormalBuffer,"
  EOL"                                __global float4* theLightSourceBuffer,"
  EOL"                                __global float4* theMaterialBuffer,"
  EOL"                                const float4 theGlobalAmbient,"
  EOL"                                const int theLightBufferSize,"
  EOL"                                const int theShadowsEnabled,"
  EOL"                                const int theReflectEnabled,"
  EOL"                                const float theSceneEpsilon,"
  EOL"                                const float theSceneRadius)"
  EOL"  {"
  EOL"    const int aPixelX = get_global_id (0);"
  EOL"    const int aPixelY = get_global_id (1);"
  EOL
  EOL"    if (aPixelX >= theSizeX || aPixelY >= theSizeY)"
  EOL"      return;"
  EOL
  EOL"    private SRay aRay;"
  EOL
  EOL"    GenerateRay (&aRay,"
  EOL"                 aPixelX,"
  EOL"                 aPixelY,"
  EOL"                 theSizeX,"
  EOL"                 theSizeY,"
  EOL"                 theOrigins,"
  EOL"                 theDirects);"
  EOL
  EOL"    float4 aColor = (float4) (0.f, 0.f, 0.f, 1.f);"
  EOL
  EOL"    float aTimeStart = IntersectBox (&aRay, theSceneMinPointBuffer[0], theSceneMaxPointBuffer[0]);"
  EOL
  EOL"    if (aTimeStart != MAXFLOAT)"
  EOL"    {"
  EOL"      aRay.Origin += aRay.Direct * max (aTimeStart - theSceneEpsilon, 0.f);"
  EOL
  EOL"      aColor = clamp (Radiance (&aRay,"
  EOL"                                theEnvMap,"
  EOL"                                theSceneNodeInfoBuffer,"
  EOL"                                theSceneMinPointBuffer,"
  EOL"                                theSceneMaxPointBuffer,"
  EOL"                                theObjectNodeInfoBuffer,"
  EOL"                                theObjectMinPointBuffer,"
  EOL"                                theObjectMaxPointBuffer,"
  EOL"                                theGeometryTriangBuffer,"
  EOL"                                theGeometryVertexBuffer,"
  EOL"                                theGeometryNormalBuffer,"
  EOL"                                theLightSourceBuffer,"
  EOL"                                theMaterialBuffer,"
  EOL"                                theGlobalAmbient,"
  EOL"                                theLightBufferSize,"
  EOL"                                theShadowsEnabled,"
  EOL"                                theReflectEnabled,"
  EOL"                                theSceneEpsilon,"
  EOL"                                theSceneRadius), 0.f, 1.f);"
  EOL"    }"
  EOL
  EOL"    write_imagef (theOutput, (int2) (aPixelX, aPixelY), aColor);"
  EOL"  }"
  EOL
  EOL"  const sampler_t OutputSampler ="
  EOL"            CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;"
  EOL
  EOL"  #define _LUM_DELTA_ 0.085f"
  EOL
  EOL"  #define AA_MAX 0.559017f"
  EOL"  #define AA_MIN 0.186339f"
  EOL
  // =======================================================================
  // function : RaytraceSmooth
  // purpose  : Performs adaptive sub-pixel rendering
  // =======================================================================
  EOL"  __kernel void RaytraceSmooth (const int theSizeX,"
  EOL"                                const int theSizeY,"
  EOL"                                const float16 theOrigins,"
  EOL"                                const float16 theDirects,"
  EOL"                                __read_only image2d_t theInput,"
  EOL"                                __read_only image2d_t theEnvMap,"
  EOL"                                __write_only image2d_t theOutput,"
  EOL"                                __global int4* theSceneNodeInfoBuffer,"
  EOL"                                __global float4* theSceneMinPointBuffer,"
  EOL"                                __global float4* theSceneMaxPointBuffer,"
  EOL"                                __global int4* theObjectNodeInfoBuffer,"
  EOL"                                __global float4* theObjectMinPointBuffer,"
  EOL"                                __global float4* theObjectMaxPointBuffer,"
  EOL"                                __global int4* theGeometryTriangBuffer,"
  EOL"                                __global float4* theGeometryVertexBuffer,"
  EOL"                                __global float4* theGeometryNormalBuffer,"
  EOL"                                __global float4* theLightSourceBuffer,"
  EOL"                                __global float4* theMaterialBuffer,"
  EOL"                                const float4 theGlobalAmbient,"
  EOL"                                const int theLightBufferSize,"
  EOL"                                const int theShadowsEnabled,"
  EOL"                                const int theReflectEnabled,"
  EOL"                                const float theSceneEpsilon,"
  EOL"                                const float theSceneRadius)"
  EOL"  {"
  EOL"    const int aPixelX = get_global_id (0);"
  EOL"    const int aPixelY = get_global_id (1);"
  EOL
  EOL"    if (aPixelX >= theSizeX || aPixelY >= theSizeY)"
  EOL"      return;"
  EOL
  EOL"    float4 aClr0 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 0, aPixelY + 0));"
  EOL"    float4 aClr1 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 0, aPixelY - 1));"
  EOL"    float4 aClr2 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 0, aPixelY + 1));"
  EOL
  EOL"    float4 aClr3 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 1, aPixelY + 0));"
  EOL"    float4 aClr4 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 1, aPixelY - 1));"
  EOL"    float4 aClr5 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 1, aPixelY + 1));"
  EOL
  EOL"    float4 aClr6 = read_imagef (theInput, OutputSampler, (float2) (aPixelX - 1, aPixelY + 0));"
  EOL"    float4 aClr7 = read_imagef (theInput, OutputSampler, (float2) (aPixelX - 1, aPixelY - 1));"
  EOL"    float4 aClr8 = read_imagef (theInput, OutputSampler, (float2) (aPixelX - 1, aPixelY + 1));"
  EOL
  EOL"    bool render = fabs (aClr1.w - aClr0.w) > _LUM_DELTA_ ||"
  EOL"                  fabs (aClr2.w - aClr0.w) > _LUM_DELTA_ ||"
  EOL"                  fabs (aClr3.w - aClr0.w) > _LUM_DELTA_ ||"
  EOL"                  fabs (aClr4.w - aClr0.w) > _LUM_DELTA_ ||"
  EOL"                  fabs (aClr5.w - aClr0.w) > _LUM_DELTA_ ||"
  EOL"                  fabs (aClr6.w - aClr0.w) > _LUM_DELTA_ ||"
  EOL"                  fabs (aClr7.w - aClr0.w) > _LUM_DELTA_ ||"
  EOL"                  fabs (aClr8.w - aClr0.w) > _LUM_DELTA_;"
  EOL
  EOL"    if (!render)"
  EOL"    {"
  EOL"      aClr1 = (aClr1.w == 1.f) ? -UNIT : aClr1;"
  EOL"      aClr2 = (aClr2.w == 1.f) ? -UNIT : aClr2;"
  EOL"      aClr3 = (aClr3.w == 1.f) ? -UNIT : aClr3;"
  EOL"      aClr4 = (aClr4.w == 1.f) ? -UNIT : aClr4;"
  EOL"      aClr5 = (aClr5.w == 1.f) ? -UNIT : aClr5;"
  EOL"      aClr6 = (aClr6.w == 1.f) ? -UNIT : aClr6;"
  EOL"      aClr7 = (aClr7.w == 1.f) ? -UNIT : aClr7;"
  EOL"      aClr8 = (aClr8.w == 1.f) ? -UNIT : aClr8;"
  EOL
  EOL"      float aLum = (aClr0.w == 1.f) ? -1.f : (0.2126f * aClr0.x + 0.7152f * aClr0.y + 0.0722f * aClr0.z);"
  EOL
  EOL"      render = fabs (0.2126f * aClr1.x + 0.7152f * aClr1.y + 0.0722f * aClr1.z - aLum) > _LUM_DELTA_ ||"
  EOL"               fabs (0.2126f * aClr2.x + 0.7152f * aClr2.y + 0.0722f * aClr2.z - aLum) > _LUM_DELTA_ ||"
  EOL"               fabs (0.2126f * aClr3.x + 0.7152f * aClr3.y + 0.0722f * aClr3.z - aLum) > _LUM_DELTA_ ||"
  EOL"               fabs (0.2126f * aClr4.x + 0.7152f * aClr4.y + 0.0722f * aClr4.z - aLum) > _LUM_DELTA_ ||"
  EOL"               fabs (0.2126f * aClr5.x + 0.7152f * aClr5.y + 0.0722f * aClr5.z - aLum) > _LUM_DELTA_ ||"
  EOL"               fabs (0.2126f * aClr6.x + 0.7152f * aClr6.y + 0.0722f * aClr6.z - aLum) > _LUM_DELTA_ ||"
  EOL"               fabs (0.2126f * aClr7.x + 0.7152f * aClr7.y + 0.0722f * aClr7.z - aLum) > _LUM_DELTA_ ||"
  EOL"               fabs (0.2126f * aClr8.x + 0.7152f * aClr8.y + 0.0722f * aClr8.z - aLum) > _LUM_DELTA_;"
  EOL"    }"
  EOL
  EOL"    float4 aColor = clamp (aClr0, 0.f, 1.f);"
  EOL
  EOL"    private SRay aRay;"
  EOL
  EOL"    const float4 aBoxMin = theSceneMinPointBuffer[0];"
  EOL"    const float4 aBoxMax = theSceneMaxPointBuffer[0];"
  EOL
  EOL"    if (render)"
  EOL"    {"
  EOL"      for (int aSample = 0; aSample <= 3; ++aSample)"
  EOL"      {"
  EOL"          float fX = aPixelX, fY = aPixelY;"
  EOL
  EOL"          if (aSample == 0)"
  EOL"          {"
  EOL"            fX -= AA_MIN; fY -= AA_MAX;"
  EOL"          }"
  EOL"          else if (aSample == 1)"
  EOL"          {"
  EOL"            fX -= AA_MAX; fY += AA_MIN;"
  EOL"          }"
  EOL"          else if (aSample == 2)"
  EOL"          {"
  EOL"            fX += AA_MIN; fY += AA_MAX;"
  EOL"          }"
  EOL"          else"
  EOL"          {"
  EOL"            fX += AA_MAX; fY -= AA_MIN;"
  EOL"          }"
  EOL
  EOL"          GenerateRay (&aRay,"
  EOL"                       fX,"
  EOL"                       fY,"
  EOL"                       theSizeX,"
  EOL"                       theSizeY,"
  EOL"                       theOrigins,"
  EOL"                       theDirects);"
  EOL
  EOL"          float aTimeStart = IntersectBox (&aRay, aBoxMin, aBoxMax);"
  EOL
  EOL"          if (aTimeStart != MAXFLOAT)"
  EOL"          {"
  EOL"            aRay.Origin += aRay.Direct * max (aTimeStart - theSceneEpsilon, 0.f);"
  EOL
  EOL"            aColor += clamp (Radiance (&aRay,"
  EOL"                                       theEnvMap,"
  EOL"                                       theSceneNodeInfoBuffer,"
  EOL"                                       theSceneMinPointBuffer,"
  EOL"                                       theSceneMaxPointBuffer,"
  EOL"                                       theObjectNodeInfoBuffer,"
  EOL"                                       theObjectMinPointBuffer,"
  EOL"                                       theObjectMaxPointBuffer,"
  EOL"                                       theGeometryTriangBuffer,"
  EOL"                                       theGeometryVertexBuffer,"
  EOL"                                       theGeometryNormalBuffer,"
  EOL"                                       theLightSourceBuffer,"
  EOL"                                       theMaterialBuffer,"
  EOL"                                       theGlobalAmbient,"
  EOL"                                       theLightBufferSize,"
  EOL"                                       theShadowsEnabled,"
  EOL"                                       theReflectEnabled,"
  EOL"                                       theSceneEpsilon,"
  EOL"                                       theSceneRadius), 0.f, 1.f);"
  EOL"          }"
  EOL"          else"
  EOL"            aColor += (float4) (0.f, 0.f, 0.f, 1.f);"
  EOL"        }"
  EOL
  EOL"        aColor *= 1.f / 5.f;"
  EOL"    }"
  EOL
  EOL"    write_imagef (theOutput, (int2) (aPixelX, aPixelY), aColor);"
  EOL"  }";

#endif