File indexing completed on 2026-04-09 07:48:55
0001 #pragma once
0002
0003 #if defined(__CUDACC__) || defined(__CUDABE__)
0004 # define TRANCHE_FUNC __forceinline__ __device__
0005 #else
0006 # define TRANCHE_FUNC inline
0007 #endif
0008
0009
0010
0011
0012
0013
0014
0015
0016
0017
0018
0019
0020
0021
0022
0023
0024
0025
0026
0027
0028
0029
0030 #define TRANCHE_STACK_SIZE 4
0031
0032
0033 struct Tranche
0034 {
0035 float tmin[TRANCHE_STACK_SIZE] ;
0036 unsigned slice[TRANCHE_STACK_SIZE] ;
0037 int curr ;
0038 };
0039
0040
0041 TRANCHE_FUNC
0042 int tranche_push(Tranche& tr, const unsigned slice, const float tmin)
0043 {
0044 if(tr.curr >= TRANCHE_STACK_SIZE - 1) return ERROR_TRANCHE_OVERFLOW ;
0045 tr.curr++ ;
0046 tr.slice[tr.curr] = slice ;
0047 tr.tmin[tr.curr] = tmin ;
0048 return 0 ;
0049 }
0050
0051 TRANCHE_FUNC
0052 int tranche_pop(Tranche& tr, unsigned& slice, float& tmin)
0053 {
0054 if(tr.curr < 0) return ERROR_POP_EMPTY ;
0055 slice = tr.slice[tr.curr] ;
0056 tmin = tr.tmin[tr.curr] ;
0057 tr.curr-- ;
0058 return 0 ;
0059 }
0060
0061 TRANCHE_FUNC
0062 unsigned long long tranche_repr(Tranche& tr)
0063 {
0064 unsigned long long val = 0ull ;
0065 if(tr.curr == -1) return val ;
0066
0067 unsigned long long c = tr.curr ;
0068 val |= ((c+1ull)&0xfull) ;
0069
0070 do {
0071 unsigned long long x = tr.slice[c] & 0xffff ;
0072 val |= x << ((4ull-c-1ull)*16ull) ;
0073 }
0074 while(c--) ;
0075
0076 return val ;
0077 }
0078
0079
0080