eic-opticks/CSG/f4_stack.h

0001 #pragma once
0002 /**
0003 f4_stack.h : struct F4_Stack using float4 as micro stack
0004 =============================================================
0005
0006 **/
0007 #if defined(__CUDACC__) || defined(__CUDABE__)
0008 #    define F4_FUNC __forceinline__ __device__ __host__
0009 #else
0010 #    define F4_FUNC inline
0011 #endif
0012
0013
0014
0015 struct F4_Stack
0016 {
0017     float4 data  ;
0018     int curr ;
0019
0020     int push(float f)
0021     {
0022         curr += 1 ;    // curr must start at -1
0023         switch(curr)
0024         {
0025             case 0:  data.x = f ;
0026             case 1:  data.y = f ;
0027             case 2:  data.z = f ;
0028             case 3:  data.w = f ;
0029         }
0030         return 0 ;
0031     }
0032     int pop(float& f0 )
0033     {
0034         switch(curr)
0035         {
0036             case 0: f0 = data.x ; break ;
0037             case 1: f0 = data.y ; break ;
0038             case 2: f0 = data.z ; break ;
0039             case 3: f0 = data.w ; break ;
0040         }
0041         curr -= 1  ;
0042         return 0 ;
0043     }
0044
0045     int push2(float f0, float f1)
0046     {
0047         switch(curr)
0048         {
0049             case -1: data.x = f0 ; data.y = f1 ; break ;
0050             case  0: data.y = f0 ; data.z = f1 ; break ;
0051             case  1: data.z = f0 ; data.w = f1 ; break ;
0052         }
0053         curr += 2 ;    // curr must start at -1
0054         return 0 ;
0055     }
0056
0057     int pop2(float& f0, float& f1)
0058     {
0059         switch(curr)
0060         {
0061            case 0: f0 = -1.f   ; f1 = -1.f   ; break ;
0062            case 1: f0 = data.x ; f1 = data.y ; break ;
0063            case 2: f0 = data.y ; f1 = data.z ; break ;
0064            case 3: f0 = data.z ; f1 = data.w ; break ;
0065         }
0066         curr -= 2  ;
0067         return 0 ;
0068     }
0069 };
0070
0071
0072