PULSAR  2.0.0
Parallel Ultra-Light Systolic Array Runtime
 All Data Structures Files Functions Typedefs Enumerations Macros Groups
prt_vdp.c
Go to the documentation of this file.
1 
11 #include "prt_vdp.h"
12 
14 
29  int *tuple, int counter,
30  prt_vdp_function_t function,
31  size_t local_store_size,
32  int num_inputs, int num_outputs, int color)
33 {
34  // Check input parameters.
35  prt_assert(tuple != NULL, "NULL tuple");
36  prt_assert(counter > 0, "counter not larger than zero");
37  prt_assert(function != NULL, "NULL VDP function");
38  prt_assert(num_inputs >= 0, "negative number of inputs");
39  prt_assert(num_outputs >= 0, "negative number of outputs");
40 
41  // Allocate the VDP.
42  prt_vdp_t *vdp = (prt_vdp_t*)malloc(sizeof(prt_vdp_t));
43  prt_assert(vdp != NULL, "malloc failed");
44 
45  // Initialize the VDP.
46  vdp->vsa = NULL;
47  vdp->thread = NULL;
48  vdp->device = NULL;
49  vdp->tuple = tuple;
50  vdp->counter = counter;
51  vdp->function = function;
52  vdp->color = color;
53 
54  // Initialize input channels.
55  vdp->num_inputs = num_inputs;
56  if (vdp->num_inputs > 0) {
57  vdp->input =
58  (prt_channel_t**)calloc(vdp->num_inputs, sizeof(prt_channel_t*));
59  prt_assert(vdp->input != NULL, "malloc failed");
60  }
61  // Initialize output channels.
62  vdp->num_outputs = num_outputs;
63  if (vdp->num_outputs > 0) {
64  vdp->output =
65  (prt_channel_t**)calloc(vdp->num_outputs, sizeof(prt_channel_t*));
66  prt_assert(vdp->output != NULL, "malloc failed");
67  }
68  // Allocate local store.
69  if (local_store_size > 0) {
70  vdp->local_store = (void*)malloc(local_store_size);
71  prt_assert(vdp->local_store != NULL, "malloc failed");
72  }
73  // Return the VDP.
74  return vdp;
75 }
76 
78 
88 {
89  // Check for a NULL VDP.
90  prt_assert(vdp != NULL, "NULL VDP");
91 
92  // Delete the tuple.
93  prt_assert(vdp->tuple != NULL, "NULL tuple");
94  prt_tuple_delete(vdp->tuple);
95 
96  int i;
97  // Delete all input channels.
98  for (i = 0; i < vdp->num_inputs; i++) {
99  prt_channel_t *channel = vdp->input[i];
100  if (channel != NULL)
101  prt_channel_delete(channel);
102  }
103  // Delete dangling output channels.
104  for (i = 0; i < vdp->num_outputs; i++) {
105  prt_channel_t *channel = vdp->output[i];
106  if (channel != NULL)
107  // IF a dangling thread channel.
108  if (vdp->location == PRT_LOCATION_HOST &&
109  channel->dst_node != vdp->thread->vsa->node_rank)
110  // Delete the channel.
111  prt_channel_delete(channel);
112  // ELSE IF a dangling device channel.
113  else if (vdp->location == PRT_LOCATION_DEVICE &&
114  channel->dst_node != vdp->device->vsa->node_rank)
115  // Delete the channel.
116  prt_channel_delete(channel);
117  }
118  // Free array of inputs.
119  if (vdp->num_inputs > 0)
120  free(vdp->input);
121 
122  // Free array of outputs.
123  if (vdp->num_outputs > 0)
124  free(vdp->output);
125 
126  // Free local store.
127  if (vdp->local_store != NULL)
128  free(vdp->local_store);
129 
130  if (vdp->location == PRT_LOCATION_DEVICE) {
131  cudaError_t error;
132  // Set the VDP's device.
133  error = cudaSetDevice(vdp->device->rank);
134  prt_assert(error == cudaSuccess, cudaGetErrorString(error));
135  // Destroy the VDP's stream.
136  error = cudaStreamDestroy(vdp->stream);
137  prt_assert(error == cudaSuccess, cudaGetErrorString(error));
138  }
139  // Free the VDP.
140  free(vdp);
141 }
142 
144 
153 {
154  // Check for a NULL VDP.
155  prt_assert(vdp != NULL, "NULL VDP");
156 
157  // Delete the tuple.
158  prt_assert(vdp->tuple != NULL, "NULL tuple");
159  prt_tuple_delete(vdp->tuple);
160 
161  int i;
162  // Delete all input channels.
163  for (i = 0; i < vdp->num_inputs; i++) {
164  prt_channel_t *channel = vdp->input[i];
165  if (channel != NULL)
166  prt_channel_delete(channel);
167  }
168  // Delete all output channels.
169  for (i = 0; i < vdp->num_outputs; i++) {
170  prt_channel_t *channel = vdp->output[i];
171  if (channel != NULL)
172  prt_channel_delete(channel);
173  }
174  // Free array of inputs.
175  if (vdp->num_inputs > 0)
176  free(vdp->input);
177 
178  // Free array of outputs.
179  if (vdp->num_outputs > 0)
180  free(vdp->output);
181 
182  // Free local store.
183  if (vdp->local_store != NULL)
184  free(vdp->local_store);
185 
186  // Free the VDP.
187  free(vdp);
188 }
189 
191 
201  prt_vdp_t *vdp, prt_channel_t *channel,
202  prt_channel_direction_t direction, int slot)
203 {
204  // Check input parameters.
205  prt_assert(vdp != NULL, "NULL VDP");
206  prt_assert(channel != NULL, "NULL channel");
207  prt_assert(direction == PRT_INPUT_CHANNEL || direction == PRT_OUTPUT_CHANNEL,
208  "undefined direction");
209  prt_assert(slot >= 0, "negative slot");
210 
211  // IF inserting input channel.
212  if (direction == PRT_INPUT_CHANNEL) {
213  // Check if channel destination tuple equals VDP's tuple.
214  prt_assert(prt_tuple_equal(channel->dst_tuple, vdp->tuple),
215  "input channel destination tuple does not match VDP tuple");
216  // Check if channel slot empty.
217  prt_assert(vdp->input[slot] == NULL,
218  "inserting channel in occupied input slot");
219  // Check the slot.
220  // Link the channel.
221  prt_assert(slot >= 0 && slot < vdp->num_inputs, "slot out of range");
222  vdp->input[slot] = channel;
223  // Set the destination VDP.
224  channel->dst_vdp = vdp;
225  }
226  // ELSE IF inserting output channel.
227  else if (direction == PRT_OUTPUT_CHANNEL) {
228  // Check if channel source tuple equals VDP's tuple.
229  prt_assert(prt_tuple_equal(channel->src_tuple, vdp->tuple),
230  "output channel source tuple does not match VDP tuple");
231  // Check if channel slot empty.
232  prt_assert(vdp->output[slot] == NULL,
233  "inserting channel in occupied output slot");
234  // Check the slot.
235  // Link the channel.
236  prt_assert(slot >= 0 && slot < vdp->num_outputs, "slot out of range");
237  vdp->output[slot] = channel;
238  // Set the source VDP.
239  channel->src_vdp = vdp;
240  }
241 }
242 
244 
258 prt_packet_t *prt_vdp_packet_new(prt_vdp_t *vdp, size_t size, void *data)
259 {
260  // Check input parameters.
261  prt_assert(vdp != NULL, "NULL VDP");
262  prt_assert(size > 0, "packet size equals zero");
263  prt_assert(size <= INT_MAX, "packet size larger than INT_MAX");
264 
265  // Test for NULL VDP.
266  prt_assert(vdp != NULL, "NULL VDP");
267 
268  // IF host VDP.
269  if (vdp->location == PRT_LOCATION_HOST)
270  // Create host packet.
271  return prt_packet_new_host(size, data);
272 
273  // IF device VDP.
274  if (vdp->location == PRT_LOCATION_DEVICE)
275  // Create device packet.
276  return prt_packet_new_device(size, data, vdp);
277 
278  prt_error("NULL thread and device");
279 }
280 
282 
298  prt_vdp_t *vdp, size_t size, void *data)
299 {
300  // Check input parameters.
301  prt_assert(vdp != NULL, "NULL VDP");
302  prt_assert(size > 0, "packet size equals zero");
303  prt_assert(size <= INT_MAX, "packet size larger than INT_MAX");
304  prt_assert(data != NULL, "NULL data pointer");
305 
306  // Create a new device packet.
307  prt_packet_t *packet = prt_packet_new_device(size, NULL, vdp);
308 
309  // Put host to device copy in the VDP's stream.
310  svg_trace_start_dma(vdp->stream);
311  cudaMemcpyAsync(
312  packet->data, data, size, cudaMemcpyHostToDevice, vdp->stream);
313  svg_trace_stop_dma(vdp->stream, Silver);
314 
315  // Return the packet.
316  return packet;
317 }
318 
320 
331 {
332  // Check input parameters.
333  prt_assert(vdp != NULL, "NULL VDP");
334  prt_assert(packet != NULL, "NULL packet");
335 
336  // IF a host packet.
337  if (vdp->location == PRT_LOCATION_HOST) {
338  // Release immediately.
339  prt_packet_release_host(packet);
340  }
341  else {
342  // Put a callback in the VDP's stream.
343  prt_callback_release_t *callback =
344  prt_callback_release_new(vdp, packet);
345  __sync_fetch_and_add(&vdp->vsa->proxy->num_callbacks, 1);
346  cudaStreamAddCallback(
347  vdp->stream, prt_callback_release_handler, (void*)callback, 0);
348  }
349 }
350 
352 
360 void prt_vdp_channel_push(prt_vdp_t *vdp, int channel_num, prt_packet_t *packet)
361 {
362  // Check input params.
363  prt_assert(vdp != NULL, "NULL VDP");
364  prt_assert(channel_num >= 0 && channel_num < vdp->num_outputs,
365  "channel number out of range");
366  prt_assert(packet != NULL, "NULL packet");
367  prt_assert(packet->size <= vdp->output[channel_num]->size,
368  "packet size larger than channel size");
369 
370  // IF thread VDP.
371  if (vdp->location == PRT_LOCATION_HOST)
372  // Push from host.
373  prt_channel_push_host(vdp, vdp->output[channel_num], packet);
374  // ELSE IF device VPD.
375  else if (vdp->location == PRT_LOCATION_DEVICE)
376  // Push from device.
377  prt_channel_push_device(vdp, vdp->output[channel_num], packet);
378  // ELSE report error.
379  else
380  prt_error("NULL thread and device");
381 }
382 
384 
394 {
395  // Check input parameters.
396  prt_assert(vdp != NULL, "NULL VDP");
397  prt_assert(channel_num >= 0 && channel_num < vdp->num_inputs,
398  "channel number out of range");
399  prt_assert(vdp->input[channel_num]->active,
400  "popping from an inactive channel");
401 
402  // Call the channel pop function.
403  return prt_channel_pop(vdp->input[channel_num]);
404 }
405 
407 
414 void prt_vdp_channel_off(prt_vdp_t *vdp, int channel_num)
415 {
416  // Check input parameters.
417  prt_assert(vdp != NULL, "NULL VDP");
418  prt_assert(channel_num >= 0 && channel_num < vdp->num_inputs,
419  "channel number out of range");
420 
421  // Switch off the channel.
422  prt_channel_off(vdp->input[channel_num]);
423 }
424 
426 
433 void prt_vdp_channel_on(prt_vdp_t *vdp, int channel_num)
434 {
435  // Check input parameters.
436  prt_assert(vdp != NULL, "NULL VDP");
437  prt_assert(channel_num >= 0 && channel_num < vdp->num_inputs,
438  "channel number out of range");
439 
440  // Switch in the channel.
441  prt_channel_on(vdp->input[channel_num]);
442 }
443 
445 
457 {
458  // Check input parameters.
459  prt_assert(vdp != NULL, "NULL VDP");
460 
461  int i;
462  // FOR each input channel.
463  for (i = 0; i < vdp->num_inputs; i++)
464  // IF the channel was established
465  // AND the channel is active
466  // AND the list of packets is empty.
467  if (vdp->input[i] != NULL &&
468  vdp->input[i]->active &&
469  prt_channel_empty(vdp->input[i]))
470  // Return "not ready".
471  return 0;
472 
473  // Return "ready".
474  return 1;
475 }