void
EventView::selectionChanged(QListViewItem * _item)
{
if (!internalSetSelection_) {
ACE_Time_Value t = eventTime(_item);
fileSet_->coursorTimeRel(t);
}
}
MSBoolean MSDoubleClick::isDoubleClick(const XEvent *ev_)
{
if (ev_->type==ButtonPress||ev_->type==ButtonRelease)
{
if (ev_->xbutton.time-lastEventTime()<=applicationDoubleClickInterval())
{
eventTime(0);
return MSTrue;
}
else
{
eventTime(ev_->xbutton.time);
return MSFalse;
}
}
else return MSFalse;
}
void
EventView::setHistory(unsigned int _size)
{
MIRO_ASSERT(_size > 0);
if (_size > history_) {
history_ = _size;
#ifdef LSB_Q3LISTVIEW
const QList<QTreeWidgetItem*> selectedItems = list_->selectedItems();
assert(!selectedItems.isEmpty());
QTreeWidgetItem * const pSelectedTreeWidgetItem = selectedItems[0];
fileSet_->getEvents(eventTime(pSelectedTreeWidgetItem), history_);
#else
fileSet_->getEvents(eventTime(list_->selectedItem()), history_);
#endif
}
else {
history_ = _size;
pruneHistory();
}
}
void
EventView::setHistory(unsigned int _size)
{
MIRO_ASSERT(_size > 0);
if (_size > history_) {
history_ = _size;
fileSet_->getEvents(eventTime(list_->selectedItem()), history_);
}
else {
history_ = _size;
pruneHistory();
}
}
Animation::Animation(models::Animation model) {
m_slide = 0;
m_lastUpdate = eventTime();
//for now, just use a universal interval
if (model.Images.size()) {
m_interval = model.Images[0].Interval;
}
for (auto img : model.Images) {
add(checkoutImage(img.Image));
}
m_key = toJson(model);
}
/*------------------------------------------------------
** ForwardSub() -- Forward substitution of Gaussian
** elimination.
**------------------------------------------------------
*/
void ForwardSub(cl_context context, float *a, float *b, float *m, int size,int timing){
// 1. set up kernels
cl_kernel fan1_kernel,fan2_kernel;
cl_int status=0;
cl_program gaussianElim_program;
cl_event writeEvent,kernelEvent,readEvent;
float writeTime=0,readTime=0,kernelTime=0;
float writeMB=0,readMB=0;
gaussianElim_program = cl_compileProgram(
(char *)"gaussianElim_kernels.cl",NULL);
fan1_kernel = clCreateKernel(
gaussianElim_program, "Fan1", &status);
status = cl_errChk(status, (char *)"Error Creating Fan1 kernel",true);
if(status)exit(1);
fan2_kernel = clCreateKernel(
gaussianElim_program, "Fan2", &status);
status = cl_errChk(status, (char *)"Error Creating Fan2 kernel",true);
if(status)exit(1);
// 2. set up memory on device and send ipts data to device
cl_mem a_dev, b_dev, m_dev;
cl_int error=0;
a_dev = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(float)*size*size, NULL, &error);
b_dev = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(float)*size, NULL, &error);
m_dev = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(float) * size * size, NULL, &error);
command_queue = cl_getCommandQueue();
error = clEnqueueWriteBuffer(command_queue,
a_dev,
1, // change to 0 for nonblocking write
0, // offset
sizeof(float)*size*size,
a,
0,
NULL,
&writeEvent);
if (timing) writeTime+=eventTime(writeEvent,command_queue);
clReleaseEvent(writeEvent);
error = clEnqueueWriteBuffer(command_queue,
b_dev,
1, // change to 0 for nonblocking write
0, // offset
sizeof(float)*size,
b,
0,
NULL,
&writeEvent);
if (timing) writeTime+=eventTime(writeEvent,command_queue);
clReleaseEvent(writeEvent);
error = clEnqueueWriteBuffer(command_queue,
m_dev,
1, // change to 0 for nonblocking write
0, // offset
sizeof(float)*size*size,
m,
0,
NULL,
&writeEvent);
if (timing) writeTime+=eventTime(writeEvent,command_queue);
clReleaseEvent(writeEvent);
writeMB = (float)(sizeof(float) * size * (size + size + 1) / 1e6);
// 3. Determine block sizes
size_t globalWorksizeFan1[1];
size_t globalWorksizeFan2[2];
size_t localWorksizeFan1Buf[1]={BLOCK_SIZE_0};
size_t localWorksizeFan2Buf[2]={BLOCK_SIZE_1_X, BLOCK_SIZE_1_Y};
size_t *localWorksizeFan1=NULL;
size_t *localWorksizeFan2=NULL;
globalWorksizeFan1[0] = size;
globalWorksizeFan2[0] = size;
globalWorksizeFan2[1] = size;
if(localWorksizeFan1Buf[0]){
localWorksizeFan1=localWorksizeFan1Buf;
globalWorksizeFan1[0]=(int)ceil(globalWorksizeFan1[0]/(double)localWorksizeFan1Buf[0])*localWorksizeFan1Buf[0];
}
if(localWorksizeFan2Buf[0]){
localWorksizeFan2=localWorksizeFan2Buf;
globalWorksizeFan2[0]=(int)ceil(globalWorksizeFan2[0]/(double)localWorksizeFan2Buf[0])*localWorksizeFan2Buf[0];
globalWorksizeFan2[1]=(int)ceil(globalWorksizeFan2[1]/(double)localWorksizeFan2Buf[1])*localWorksizeFan2Buf[1];
}
int t;
// 4. Setup and Run kernels
for (t=0; t<(size-1); t++) {
// kernel args
cl_int argchk;
argchk = clSetKernelArg(fan1_kernel, 0, sizeof(cl_mem), (void *)&m_dev);
argchk |= clSetKernelArg(fan1_kernel, 1, sizeof(cl_mem), (void *)&a_dev);
argchk |= clSetKernelArg(fan1_kernel, 2, sizeof(cl_mem), (void *)&b_dev);
argchk |= clSetKernelArg(fan1_kernel, 3, sizeof(int), (void *)&size);
argchk |= clSetKernelArg(fan1_kernel, 4, sizeof(int), (void *)&t);
cl_errChk(argchk,"ERROR in Setting Fan1 kernel args",true);
//printf("localWorksizeFan1:%u, globalWorksizeFan1:%u\n", localWorksizeFan1Buf[0], globalWorksizeFan1[0]);
#pragma dividend local_work_group_size localWorksizeFan1 dim 1 dim1(2:64:2:64)
//This lws will be used to profile the OpenCL kernel with id 1
size_t _dividend_lws_localWorksizeFan1_k1[2];
{
_dividend_lws_localWorksizeFan1_k1[0] = getLWSValue("DIVIDEND_LWS1_D0",DIVIDEND_LWS1_D0_DEFAULT_VAL);
//Dividend extension: store the kernel id as the last element
_dividend_lws_localWorksizeFan1_k1[1] = 1;
}
// launch kernel
error = DIVIDEND_CL_WRAP(clEnqueueNDRangeKernel)(
command_queue, fan1_kernel, 1, 0,
globalWorksizeFan1, _dividend_lws_localWorksizeFan1_k1,
0, NULL, NULL);
cl_errChk(error,"ERROR in Executing Fan1 Kernel",true);
//fprintf(stderr, "AFTER THIS\n");
argchk = clSetKernelArg(fan2_kernel, 0, sizeof(cl_mem), (void *)&m_dev);
argchk |= clSetKernelArg(fan2_kernel, 1, sizeof(cl_mem), (void *)&a_dev);
argchk |= clSetKernelArg(fan2_kernel, 2, sizeof(cl_mem), (void *)&b_dev);
argchk |= clSetKernelArg(fan2_kernel, 3, sizeof(int), (void *)&size);
argchk |= clSetKernelArg(fan2_kernel, 4, sizeof(int), (void *)&t);
cl_errChk(argchk,"ERROR in Setting Fan2 kernel args",true);
size_t local_work_size[] = {128, 128};
//printf("localWorksizeFan2:%u, globalWorksizeFan2[0]:%u, globalWorksizeFan2[1]:%u\n", localWorksizeFan2Buf[0], globalWorksizeFan2[0], globalWorksizeFan2[1]);
#pragma dividend local_work_group_size local_work_size dim 2 dim1(8:64:2:64) dim2(8:64:2:64)
//This lws will be used to profile the OpenCL kernel with id 2
size_t _dividend_lws_local_work_size_k2[3];
{
_dividend_lws_local_work_size_k2[0] = getLWSValue("DIVIDEND_LWS2_D0",DIVIDEND_LWS2_D0_DEFAULT_VAL);
_dividend_lws_local_work_size_k2[1] = getLWSValue("DIVIDEND_LWS2_D1",DIVIDEND_LWS2_D1_DEFAULT_VAL);
//Dividend extension: store the kernel id as the last element
_dividend_lws_local_work_size_k2[2] = 2;
}
// launch kernel
error = DIVIDEND_CL_WRAP(clEnqueueNDRangeKernel)(
command_queue, fan2_kernel, 2, 0,
globalWorksizeFan2, _dividend_lws_local_work_size_k2,
0, NULL, NULL);
cl_errChk(error,"ERROR in Executing Fan2 Kernel",true);
if (timing) {
// printf("here2a\n");
// kernelTime+=eventTime(kernelEvent,command_queue);
// printf("here2b\n");
}
clReleaseEvent(kernelEvent);
//Fan2<<<dimGridXY,dimBlockXY>>>(m_cuda,a_cuda,b_cuda,Size,Size-t,t);
//cudaThreadSynchronize();
}
// 5. transfer data off of device
error = clEnqueueReadBuffer(command_queue,
a_dev,
1, // change to 0 for nonblocking write
0, // offset
sizeof(float) * size * size,
a,
0,
NULL,
&readEvent);
cl_errChk(error,"ERROR with clEnqueueReadBuffer",true);
if (timing) readTime+=eventTime(readEvent,command_queue);
clReleaseEvent(readEvent);
error = clEnqueueReadBuffer(command_queue,
b_dev,
1, // change to 0 for nonblocking write
0, // offset
sizeof(float) * size,
b,
0,
NULL,
&readEvent);
cl_errChk(error,"ERROR with clEnqueueReadBuffer",true);
if (timing) readTime+=eventTime(readEvent,command_queue);
clReleaseEvent(readEvent);
error = clEnqueueReadBuffer(command_queue,
m_dev,
1, // change to 0 for nonblocking write
0, // offset
sizeof(float) * size * size,
m,
0,
NULL,
&readEvent);
cl_errChk(error,"ERROR with clEnqueueReadBuffer",true);
if (timing) readTime+=eventTime(readEvent,command_queue);
clReleaseEvent(readEvent);
readMB = (float)(sizeof(float) * size * (size + size + 1) / 1e6);
if (timing) {
printf("Matrix Size\tWrite(s) [size]\t\tKernel(s)\tRead(s) [size]\t\tTotal(s)\n");
printf("%dx%d \t",size,size);
printf("%f [%.2fMB]\t",writeTime,writeMB);
printf("%f\t",kernelTime);
printf("%f [%.2fMB]\t",readTime,readMB);
printf("%f\n\n",writeTime+kernelTime+readTime);
}
}
