27 #ifndef bits_gpulog_msg_layout_h__
28 #define bits_gpulog_msg_layout_h__
44 inline std::ostream &operator <<(std::ostream &out, const array<T> &v) {
return out <<
"n=" << v.nelem; }
45 inline std::ostream &
operator <<(std::ostream &out,
const header &h) {
return out <<
"msgid=" << h.msgid <<
" len=" << h.len IFARGINFO( <<
" nargs=" << h.nargs <<
" infos=" << h.infos); }
46 inline std::ostream &
operator <<(std::ostream &out,
const Tunspec &tu) {
return out; }
47 inline std::ostream &
operator <<(std::ostream &out,
const arginfo &h)
49 out <<
"arg=" << h.arg <<
" align=" << h.align <<
" size=" << h.size <<
" dim=" << h.dim;
50 out <<
" isarr=" << h.isarray <<
" nelem=" << h.nelem;;
51 out <<
" begin=" << h.begin <<
" len=" << h.len;
65 template<
typename T>
struct argio
67 __host__ __device__
static inline void put(
char *ptr,
const T &
x,
int start,
int datalen)
69 DHOST( std::cerr <<
"Writing [" << x <<
"] start=" << start <<
" len=" << datalen <<
"\n" );
70 DGPU( printf(
"Writing start=%d len=%d\n", start, datalen); );
71 *(T*)(ptr + start) =
x;
76 template<
typename T>
struct argio<T*>
78 __host__ __device__
static inline void put(
char *ptr,
const T *
x,
int start,
int datalen)
80 STATIC_ASSERTION_FAILED__Pointer_serialization_is_not_allowed__Use_fixed_size_C_arrays_or_array_template_instead_____(x);
85 template<
typename T,
int N>
struct argio<T[N]>
88 __host__ __device__
static inline void put(
char *ptr,
const T *
x,
int start,
int datalen)
90 DHOST( std::cerr <<
"Writing presized array [" << x <<
"] start=" << start <<
" len=" << datalen <<
"\n" );
91 DGPU( printf(
"Writing presized array start=%d len=%d\n", start, datalen); );
93 for(
int i = 0; i != datalen /
sizeof(T); i++)
103 __host__ __device__
static inline void put(
char *ptr,
const array<T> &
x,
int start,
int datalen)
106 DHOST( std::cerr <<
"Allocating array [" << x <<
"] start=" << start <<
" element_size=" << datalen <<
" v[0]= " << *(T*)(ptr + start) <<
"\n" );
107 DGPU( printf(
"Allocating array start=%d element_size=%d\n", start, datalen); );
112 #define ASTART(at, a) (at & (a-1) ? (at & ~(a-1)) + a : at)
117 template<
int at,
int a>
struct ata {
static const int value =
ASTART(at, a); };
118 #define ASTARTx(at, a) (ata<at, a>::value)
120 #define ADDR(beg, k) \
121 static const int align##k = ALIGNOF(T##k); \
122 static const int begin##k = ASTARTx(beg, align##k); \
123 static const int len##k = SIZEOF(T##k); \
124 static const int end##k = begin##k + len##k; \
125 typedef argio<T##k> IO##k; \
126 __host__ __device__ static inline void dump##k() { if(ISUNSPEC(T##k)) { return; }; DHOST( std::cerr << "arg " << k << ": begin=" << begin##k << " end=" << end##k << " len=" << len##k << " isarray=" << ISARRAY(T##k) << "\n"; ); } \
127 __host__ __device__ static inline void get_arginfo##k(arginfo *ai, int nelem) \
129 typedef ttrait<T##k> TT; \
130 if(!ISUNSPEC(T##k)) \
133 ai->align = TT::align; \
134 ai->size = TT::size; \
136 ai->isarray = TT::isarr; \
137 ai->begin = begin##k; \
138 ai->nelem = ai->isarray ? nelem : 1; \
139 ai->len = len##k * ai->nelem; \
149 typename T1 = Tunspec,
typename T2 = Tunspec,
typename T3 = Tunspec,
typename T4 = Tunspec,
typename T5 = Tunspec,
150 typename T6 = Tunspec,
typename T7 = Tunspec,
typename T8 = Tunspec,
typename T9 = Tunspec,
typename T10 = Tunspec
155 static const int nargs =
157 + !ISUNSPEC(T1) + !ISUNSPEC(T2) + !ISUNSPEC(T3) + !ISUNSPEC(T4) + !ISUNSPEC(T5) +
158 + !ISUNSPEC(T6) + !ISUNSPEC(T7) + !ISUNSPEC(T8) + !ISUNSPEC(T9) + !ISUNSPEC(T10);
178 static const int end = end10;
181 static const int len = end;
209 __host__ __device__
inline static void store_arginfo(
const char *ptr,
const T &
x)
211 return store_arginfo_aux(ptr, 1);
215 __host__ __device__
inline static void store_arginfo(
const char *ptr,
const array<T> &
x)
217 return store_arginfo_aux(ptr, x.nelem);
220 __host__ __device__
inline static void store_arginfo_aux(
const char *ptr,
int nelem)
222 DHOST( std::cerr <<
"Storing arginfo [nargs=" << nargs <<
" nelem=" << nelem <<
"]\n"; )
223 DHOST( std::cerr <<
"Unspecified args: " << ISUNSPEC(T0) <<
" " << ISUNSPEC(T1) <<
" " << ISUNSPEC(T2) <<
" " << ISUNSPEC(T3) <<
" " << ISUNSPEC(T4) <<
" " << ISUNSPEC(T5) <<
" " << ISUNSPEC(T6) <<
" " << ISUNSPEC(T7) <<
" " << ISUNSPEC(T8) <<
" " << ISUNSPEC(T9) <<
" " << ISUNSPEC(T10) <<
"\n"; )
225 arginfo *ai = (arginfo *)(ptr + begin99);
226 get_arginfo0(ai + 0, nelem);
227 get_arginfo1(ai + 1, nelem);
228 get_arginfo2(ai + 2, nelem);
229 get_arginfo3(ai + 3, nelem);
230 get_arginfo4(ai + 4, nelem);
231 get_arginfo5(ai + 5, nelem);
232 get_arginfo6(ai + 6, nelem);
233 get_arginfo7(ai + 7, nelem);
234 get_arginfo8(ai + 8, nelem);
235 get_arginfo9(ai + 9, nelem);
236 get_arginfo10(ai + 10, nelem);
239 T0 &hdr = *(T0 *)(ptr + begin0);
245 __host__ __device__
inline static void dump()
251 dump1(); dump2(); dump3(); dump4(); dump5(); dump6(); dump7(); dump8(); dump9(); dump10();
259 #endif // bits_gpulog_msg_layout_h__