1 /* Test dispatch of events to callbacks. */
11 /* Use explicit 'copyin' clauses, to work around "'firstprivate'
12 optimizations", which will cause the value at the point of call to be used
13 (*before* any potential modifications done in callbacks), as opposed to its
14 address being taken, which then later gets dereferenced (*after* any
15 modifications done in callbacks). */
16 #define COPYIN(...) copyin(__VA_ARGS__)
19 #define DEBUG_printf(...) //__builtin_printf (__VA_ARGS__)
22 static acc_prof_reg reg
;
23 static acc_prof_reg unreg
;
24 static acc_prof_lookup_func lookup
;
25 void acc_register_library (acc_prof_reg reg_
, acc_prof_reg unreg_
, acc_prof_lookup_func lookup_
)
27 DEBUG_printf ("%s\n", __FUNCTION__
);
35 static int state
= -1;
37 #define STATE_OP(state, op) \
40 typeof (state) state_o = (state); \
43 DEBUG_printf("state: %d -> %d\n", state_o, (state)); \
48 static acc_device_t acc_device_type
;
49 static int acc_device_num
;
50 static int acc_async
= acc_async_sync
;
55 acc_event_info event_info
;
56 struct tool_info
*nested
;
58 struct tool_info
*tool_info
;
60 static void cb_device_init_start (acc_prof_info
*prof_info
, acc_event_info
*event_info
, acc_api_info
*api_info
)
62 DEBUG_printf ("%s\n", __FUNCTION__
);
68 assert (tool_info
== NULL
);
69 tool_info
= (struct tool_info
*) malloc(sizeof *tool_info
);
70 assert (tool_info
!= NULL
);
71 tool_info
->nested
= NULL
;
73 assert (prof_info
->event_type
== acc_ev_device_init_start
);
74 assert (prof_info
->valid_bytes
== _ACC_PROF_INFO_VALID_BYTES
);
75 assert (prof_info
->version
== _ACC_PROF_INFO_VERSION
);
77 assert (prof_info
->device_type
== acc_device_host
);
79 assert (prof_info
->device_type
== acc_device_default
);
80 assert (prof_info
->device_number
== acc_device_num
);
81 assert (prof_info
->thread_id
== -1);
82 assert (prof_info
->async
== acc_async_sync
);
83 assert (prof_info
->async_queue
== prof_info
->async
);
84 assert (prof_info
->src_file
== NULL
);
85 assert (prof_info
->func_name
== NULL
);
86 assert (prof_info
->line_no
== -1);
87 assert (prof_info
->end_line_no
== -1);
88 assert (prof_info
->func_line_no
== -1);
89 assert (prof_info
->func_end_line_no
== -1);
91 assert (event_info
->other_event
.event_type
== prof_info
->event_type
);
92 assert (event_info
->other_event
.valid_bytes
== _ACC_OTHER_EVENT_INFO_VALID_BYTES
);
93 assert (event_info
->other_event
.parent_construct
== acc_construct_runtime_api
);
94 assert (event_info
->other_event
.implicit
== 0);
95 assert (event_info
->other_event
.tool_info
== NULL
);
97 assert (api_info
->device_api
== acc_device_api_none
);
98 assert (api_info
->valid_bytes
== _ACC_API_INFO_VALID_BYTES
);
99 assert (api_info
->device_type
== prof_info
->device_type
);
100 assert (api_info
->vendor
== -1);
101 assert (api_info
->device_handle
== NULL
);
102 assert (api_info
->context_handle
== NULL
);
103 assert (api_info
->async_handle
== NULL
);
105 tool_info
->event_info
.other_event
.event_type
= event_info
->other_event
.event_type
;
106 event_info
->other_event
.tool_info
= tool_info
;
109 static void cb_device_init_end (acc_prof_info
*prof_info
, acc_event_info
*event_info
, acc_api_info
*api_info
)
111 DEBUG_printf ("%s\n", __FUNCTION__
);
115 STATE_OP (state
, ++);
117 assert (tool_info
!= NULL
);
118 assert (tool_info
->event_info
.other_event
.event_type
== acc_ev_device_init_start
);
120 assert (prof_info
->event_type
== acc_ev_device_init_end
);
121 assert (prof_info
->valid_bytes
== _ACC_PROF_INFO_VALID_BYTES
);
122 assert (prof_info
->version
== _ACC_PROF_INFO_VERSION
);
124 assert (prof_info
->device_type
== acc_device_host
);
126 assert (prof_info
->device_type
== acc_device_default
);
127 assert (prof_info
->device_number
== acc_device_num
);
128 assert (prof_info
->thread_id
== -1);
129 assert (prof_info
->async
== acc_async_sync
);
130 assert (prof_info
->async_queue
== prof_info
->async
);
131 assert (prof_info
->src_file
== NULL
);
132 assert (prof_info
->func_name
== NULL
);
133 assert (prof_info
->line_no
== -1);
134 assert (prof_info
->end_line_no
== -1);
135 assert (prof_info
->func_line_no
== -1);
136 assert (prof_info
->func_end_line_no
== -1);
138 assert (event_info
->other_event
.event_type
== prof_info
->event_type
);
139 assert (event_info
->other_event
.valid_bytes
== _ACC_OTHER_EVENT_INFO_VALID_BYTES
);
140 assert (event_info
->other_event
.parent_construct
== acc_construct_runtime_api
);
141 assert (event_info
->other_event
.implicit
== 0);
142 assert (event_info
->other_event
.tool_info
== tool_info
);
144 assert (api_info
->device_api
== acc_device_api_none
);
145 assert (api_info
->valid_bytes
== _ACC_API_INFO_VALID_BYTES
);
146 assert (api_info
->device_type
== prof_info
->device_type
);
147 assert (api_info
->vendor
== -1);
148 assert (api_info
->device_handle
== NULL
);
149 assert (api_info
->context_handle
== NULL
);
150 assert (api_info
->async_handle
== NULL
);
156 static void cb_compute_construct_start (acc_prof_info
*prof_info
, acc_event_info
*event_info
, acc_api_info
*api_info
)
158 DEBUG_printf ("%s\n", __FUNCTION__
);
162 STATE_OP (state
, ++);
164 assert (tool_info
== NULL
);
165 tool_info
= (struct tool_info
*) malloc(sizeof *tool_info
);
166 assert (tool_info
!= NULL
);
167 tool_info
->nested
= NULL
;
169 assert (prof_info
->event_type
== acc_ev_compute_construct_start
);
170 assert (prof_info
->valid_bytes
== _ACC_PROF_INFO_VALID_BYTES
);
171 assert (prof_info
->version
== _ACC_PROF_INFO_VERSION
);
172 assert (prof_info
->device_type
== acc_device_type
);
173 assert (prof_info
->device_number
== acc_device_num
);
174 assert (prof_info
->thread_id
== -1);
175 assert (prof_info
->async
== /* TODO acc_async */ acc_async_sync
);
176 assert (prof_info
->async_queue
== prof_info
->async
);
177 assert (prof_info
->src_file
== NULL
);
178 assert (prof_info
->func_name
== NULL
);
179 assert (prof_info
->line_no
== -1);
180 assert (prof_info
->end_line_no
== -1);
181 assert (prof_info
->func_line_no
== -1);
182 assert (prof_info
->func_end_line_no
== -1);
184 assert (event_info
->other_event
.event_type
== prof_info
->event_type
);
185 assert (event_info
->other_event
.valid_bytes
== _ACC_OTHER_EVENT_INFO_VALID_BYTES
);
186 assert (event_info
->other_event
.parent_construct
== acc_construct_parallel
);
187 assert (event_info
->other_event
.implicit
== 0);
188 assert (event_info
->other_event
.tool_info
== NULL
);
190 assert (api_info
->device_api
== acc_device_api_none
);
191 assert (api_info
->valid_bytes
== _ACC_API_INFO_VALID_BYTES
);
192 assert (api_info
->device_type
== prof_info
->device_type
);
193 assert (api_info
->vendor
== -1);
194 assert (api_info
->device_handle
== NULL
);
195 assert (api_info
->context_handle
== NULL
);
196 assert (api_info
->async_handle
== NULL
);
198 tool_info
->event_info
.other_event
.event_type
= event_info
->other_event
.event_type
;
199 event_info
->other_event
.tool_info
= tool_info
;
202 static void cb_compute_construct_end (acc_prof_info
*prof_info
, acc_event_info
*event_info
, acc_api_info
*api_info
)
204 DEBUG_printf ("%s\n", __FUNCTION__
);
209 /* In an 'async' setting, this event may be triggered before actual 'async'
210 data copying has completed. Given that 'state' appears in 'COPYIN', we
211 first have to synchronize (that is, let the 'async' 'COPYIN' read the
212 current 'state' value)... */
213 if (acc_async
!= acc_async_sync
)
215 /* "We're not yet accounting for the fact that _OpenACC events may occur
216 during event processing_"; temporarily disable to avoid deadlock. */
217 unreg (acc_ev_none
, NULL
, acc_toggle_per_thread
);
218 acc_wait (acc_async
);
219 reg (acc_ev_none
, NULL
, acc_toggle_per_thread
);
221 /* ... before modifying it in the following. */
223 STATE_OP (state
, ++);
225 assert (tool_info
!= NULL
);
226 assert (tool_info
->event_info
.other_event
.event_type
== acc_ev_compute_construct_start
);
227 assert (tool_info
->nested
== NULL
);
229 assert (prof_info
->event_type
== acc_ev_compute_construct_end
);
230 assert (prof_info
->valid_bytes
== _ACC_PROF_INFO_VALID_BYTES
);
231 assert (prof_info
->version
== _ACC_PROF_INFO_VERSION
);
232 assert (prof_info
->device_type
== acc_device_type
);
233 assert (prof_info
->device_number
== acc_device_num
);
234 assert (prof_info
->thread_id
== -1);
235 if (acc_device_type
== acc_device_host
)
236 assert (prof_info
->async
== acc_async_sync
);
238 assert (prof_info
->async
== acc_async
);
239 assert (prof_info
->async_queue
== prof_info
->async
);
240 assert (prof_info
->src_file
== NULL
);
241 assert (prof_info
->func_name
== NULL
);
242 assert (prof_info
->line_no
== -1);
243 assert (prof_info
->end_line_no
== -1);
244 assert (prof_info
->func_line_no
== -1);
245 assert (prof_info
->func_end_line_no
== -1);
247 assert (event_info
->other_event
.event_type
== prof_info
->event_type
);
248 assert (event_info
->other_event
.valid_bytes
== _ACC_OTHER_EVENT_INFO_VALID_BYTES
);
249 assert (event_info
->other_event
.parent_construct
== acc_construct_parallel
);
250 assert (event_info
->other_event
.implicit
== 0);
251 assert (event_info
->other_event
.tool_info
== tool_info
);
253 if (acc_device_type
== acc_device_host
)
254 assert (api_info
->device_api
== acc_device_api_none
);
255 else if (acc_device_type
== acc_device_radeon
)
256 assert (api_info
->device_api
== acc_device_api_other
);
258 assert (api_info
->device_api
== acc_device_api_cuda
);
259 assert (api_info
->valid_bytes
== _ACC_API_INFO_VALID_BYTES
);
260 assert (api_info
->device_type
== prof_info
->device_type
);
261 assert (api_info
->vendor
== -1);
262 assert (api_info
->device_handle
== NULL
);
263 assert (api_info
->context_handle
== NULL
);
264 assert (api_info
->async_handle
== NULL
);
273 acc_register_library (acc_prof_register
, acc_prof_unregister
, acc_prof_lookup
);
275 STATE_OP (state
, = 0);
276 reg (acc_ev_device_init_start
, cb_device_init_start
, acc_reg
);
277 reg (acc_ev_device_init_end
, cb_device_init_end
, acc_reg
);
278 reg (acc_ev_compute_construct_start
, cb_compute_construct_start
, acc_reg
);
279 reg (acc_ev_compute_construct_end
, cb_compute_construct_end
, acc_reg
);
282 acc_init (acc_device_host
);
285 STATE_OP (state
, = 10);
287 acc_device_type
= acc_get_device_type ();
288 acc_device_num
= acc_get_device_num (acc_device_type
);
293 #pragma acc parallel async(acc_async) COPYIN(state) copyout(state_init)
297 acc_async
= acc_async_sync
;
299 assert (state_init
== 11);
301 assert (state
== 12);
303 STATE_OP (state
, = 90);
304 acc_shutdown (acc_device_host
);
305 assert (state
== 90);
308 STATE_OP (state
, = 100);
309 acc_init (acc_device_default
);
310 assert (state
== 102);
312 STATE_OP (state
, = 110);
314 acc_device_type
= acc_get_device_type ();
315 acc_device_num
= acc_get_device_num (acc_device_type
);
320 #pragma acc parallel async(acc_async) COPYIN(state) copyout(state_init)
324 acc_async
= acc_async_sync
;
326 assert (state_init
== 111);
328 assert (state
== 112);
330 STATE_OP (state
, = 190);
331 acc_shutdown (acc_device_default
);
332 assert (state
== 190);