static odp_pktio_t create_pktio(int iface_idx, odp_pktio_input_mode_t imode,
odp_pktio_output_mode_t omode)
{
odp_pktio_t pktio;
odp_pktio_param_t pktio_param;
const char *iface = iface_name[iface_idx];
odp_pktio_param_init(&pktio_param);
pktio_param.in_mode = imode;
pktio_param.out_mode = omode;
pktio = odp_pktio_open(iface, pool[iface_idx], &pktio_param);
if (pktio == ODP_PKTIO_INVALID)
pktio = odp_pktio_lookup(iface);
CU_ASSERT(pktio != ODP_PKTIO_INVALID);
CU_ASSERT(odp_pktio_to_u64(pktio) !=
odp_pktio_to_u64(ODP_PKTIO_INVALID));
/* Print pktio debug info and test that the odp_pktio_print() function
* is implemented. */
if (pktio != ODP_PKTIO_INVALID)
odp_pktio_print(pktio);
if (wait_for_network)
spin_wait(ODP_TIME_SEC_IN_NS / 4);
return pktio;
}
void s_barrier()
{
if (run_cpus == 1 || v->fail_safe & 3) {
return;
}
spin_wait(&barr->s_st1); /* Wait if the barrier is active */
spin_lock(&barr->s_lck); /* Get lock for barr struct */
if (--barr->s_count == 0) { /* Last process? */
barr->s_st1.slock = 0; /* Hold up any processes re-entering */
barr->s_st2.slock = 1; /* Release the other processes */
barr->s_count++;
spin_unlock(&barr->s_lck);
} else {
spin_unlock(&barr->s_lck);
spin_wait(&barr->s_st2); /* wait for peers to arrive */
spin_lock(&barr->s_lck);
if (++barr->s_count == barr->s_maxproc) {
barr->s_st1.slock = 1;
barr->s_st2.slock = 0;
}
spin_unlock(&barr->s_lck);
}
}
void Scheduler::resume(Thread * _thread)
{
if((tag==1) && (front==rear))
{
Console::puts("Queue is full\n");
for(;;);
}
else
{
queue[rear]=_thread;
rear=(rear+1)%M;
tag=1;
spin_wait("resume a thread onto ready queue\n",5000000);
}
}
void Scheduler::yield()
{
if((tag==0) && (front==rear))
{
Console::puts("Queue is empty\n");
//for(;;);
}
else
{
Thread * nexthread=queue[front];
front=(front+1)%M;
tag=0;
spin_wait("A new thread taken out\n",5000000);
Thread::dispatch_to(nexthread);
}
}
static void unit_work(MemtestContext *ctx) {
#ifndef NO_BARRIER
# define setup_barrier sb_spin(ctx->barrier)
#else
# define setup_barrier
#endif
cl_int err;
int idx;
bool mixed_mode = ctx->shared_info->mixed_mode;
int leader_id = ctx->shared_info->leader_id;
// BARRIER
setup_barrier;
int gpu_id = ctx->id;
size_t memsize = ctx->shared_info->memsize[gpu_id];
cl_context cl_ctx = ctx->memtest->contexts[gpu_id];
cl_command_queue queue = ctx->memtest->queues[gpu_id];
cl_kernel kernel = ctx->memtest->knl_inc[gpu_id];
long long ts_delta = ctx->shared_info->gpu_ts_profile[gpu_id];
long long rw_offset_min = LLONG_MAX;
for (idx = 0; idx < ctx->memtest->ndevices; idx++) {
if (rw_offset_min > ctx->shared_info->rw_offsets[idx])
rw_offset_min = ctx->shared_info->rw_offsets[idx];
}
long long rw_offset = ctx->shared_info->rw_offsets[gpu_id] - rw_offset_min;
cl_uchar *buffer = calloc(memsize, sizeof(cl_uchar));
if (!buffer) {
fprintf(stderr, "Out of Memory\n");
abort();
}
if (ctx->shared_info->randomize_buffer) {
unsigned int seed = (unsigned int)time(NULL);
cl_uchar* b_iter, *b_end;
b_end = buffer + memsize;
for (b_iter = buffer; b_iter < b_end; b_iter++) {
*b_iter = (cl_uchar)rand_r(&seed);
}
}
// BARRIER
setup_barrier;
cl_mem mem_arr;
mem_arr = clCreateBuffer(cl_ctx, CL_MEM_READ_WRITE, sizeof(cl_uchar) * memsize, NULL, &err);
CHECK_ERROR;
fprintf(stderr, "Created a buffer\n");
fprintf(stderr, "[%d]Starting to transfer the buffer to each device..\n", gpu_id);
// BARRIER
setup_barrier;
if (mixed_mode && gpu_id != leader_id) {
setup_barrier;
setup_barrier;
setup_barrier;
}
spin_wait(rw_offset);
clock_t st_t = clock();
cl_event h2d_ev, d2h_ev, kernel_ev;
err = clEnqueueWriteBuffer(queue, mem_arr, CL_TRUE, 0, sizeof(cl_uchar) * memsize, buffer, 0, NULL, &h2d_ev);
CHECK_ERROR;
clFinish(queue);
clWaitForEvents(1, &h2d_ev);
GpuLog_log(ctx->shared_info->logger, gpu_id, "H2DMemcpy", ts_delta, h2d_ev);
clReleaseEvent(h2d_ev);
// BARRIER
setup_barrier;
fprintf(stderr, "Transfered the buffer to each device..(%.3lf)\n", ((double)st_t / CLOCKS_PER_SEC) * 1000);
size_t gblsize[1] = {memsize};
size_t lclsize[1] = {memsize > 128? 128 : memsize};
err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &mem_arr);
CHECK_ERROR;
cl_uint _size = memsize;
err = clSetKernelArg(kernel, 1, sizeof(cl_uint), &_size);
CHECK_ERROR;
// BARRIER
setup_barrier;
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, gblsize, lclsize, 0, NULL, &kernel_ev);
CHECK_ERROR;
clFinish(queue);
clWaitForEvents(1, &kernel_ev);
GpuLog_log(ctx->shared_info->logger, gpu_id, "kernel_inc", ts_delta, kernel_ev);
clReleaseEvent(kernel_ev);
// BARRIER
setup_barrier;
spin_wait(rw_offset);
err = clEnqueueReadBuffer(queue, mem_arr, CL_TRUE, 0, sizeof(cl_uchar) * memsize, buffer, 0, NULL, &d2h_ev);
CHECK_ERROR;
clFinish(queue);
clWaitForEvents(1, &d2h_ev);
GpuLog_log(ctx->shared_info->logger, gpu_id, "D2HMemcpy", ts_delta, d2h_ev);
clReleaseEvent(d2h_ev);
free(buffer);
clReleaseMemObject(mem_arr);
if (mixed_mode && gpu_id == leader_id) {
setup_barrier;
setup_barrier;
setup_barrier;
}
#undef setup_barrier
}
