eic-opticks/CSG/csg_stack.h

0001 #pragma once
0002
0003 #if defined(__CUDACC__) || defined(__CUDABE__)
0004 #    define CSG_FUNC __forceinline__ __device__ __host__
0005 #else
0006 #    define CSG_FUNC inline
0007 #endif
0008
0009
0010 #ifdef DEBUG_RECORD
0011 #include "CSGRecord.h"
0012 #endif
0013
0014
0015 /**2
0016 csg_stack.h : struct CSG_Stack, csg_push, csg_pop
0017 ----------------------------------------------------
0018
0019 Small stack of float4 isect (surface normals and t:distance_to_intersect).
0020
0021 csg_push
0022     push float4 isect and nodeIdx into stack
0023
0024 csg_pop
0025     pop float4 isect and nodeIdx off the stack
0026
0027 2**/
0028
0029 #define CSG_STACK_SIZE 15
0030
0031 struct CSG_Stack
0032 {
0033    float4 data[CSG_STACK_SIZE] ;
0034    unsigned idx[CSG_STACK_SIZE] ;
0035    int curr ;
0036 };
0037
0038 CSG_FUNC
0039 int csg_push(CSG_Stack& csg, const float4& isect, unsigned nodeIdx)
0040 {
0041 #ifdef DEBUG
0042      assert( csg.curr < CSG_STACK_SIZE );
0043 #endif
0044     if(csg.curr >= CSG_STACK_SIZE - 1) return ERROR_OVERFLOW ;
0045     csg.curr++ ;
0046     csg.data[csg.curr] = isect ;
0047     csg.idx[csg.curr] = nodeIdx ;
0048
0049     return 0 ;
0050 }
0051 CSG_FUNC
0052 int csg_pop(CSG_Stack& csg, float4& isect, unsigned& nodeIdx)
0053 {
0054     if(csg.curr < 0) return ERROR_POP_EMPTY ;
0055     isect = csg.data[csg.curr] ;
0056     nodeIdx = csg.idx[csg.curr] ;
0057     csg.idx[csg.curr] = 0u ;   // scrub the idx for debug
0058     csg.curr-- ;
0059     return 0 ;
0060 }
0061
0062 CSG_FUNC
0063 int csg_pop0(CSG_Stack& csg)   // pop without returning anything
0064 {
0065     if(csg.curr < 0) return ERROR_POP_EMPTY ;
0066     csg.idx[csg.curr] = 0u ;   // scrub the idx for debug
0067     csg.curr-- ;
0068     return 0 ;
0069 }
0070
0071 CSG_FUNC
0072 unsigned long long csg_repr(CSG_Stack& csg)
0073 {
0074     unsigned long long val = 0 ;
0075     if(csg.curr == -1) return val ;
0076
0077     unsigned long long c = csg.curr ;
0078     val |= (c+1ull) ;  // count at lsb, contents from msb
0079
0080     do {
0081         unsigned long long x = csg.idx[c] & 0xf ;
0082         val |=  x << ((16ull-c-1ull)*4ull) ;
0083     }
0084     while(c--) ;
0085
0086     return val ;
0087 }
0088
0089
0090