Skip to content
Snippets Groups Projects
Commit 6c0682bd authored by Maria Kotsifakou's avatar Maria Kotsifakou
Browse files

Edited timers in spmv

parent 8f3c406e
No related branches found
No related tags found
No related merge requests found
...@@ -77,6 +77,7 @@ int main(int argc, char** argv) { ...@@ -77,6 +77,7 @@ int main(int argc, char** argv) {
cl_kernel clKernel = clCreateKernel(clProgram,"spmv_jds",&clStatus); cl_kernel clKernel = clCreateKernel(clProgram,"spmv_jds",&clStatus);
CHECK_ERROR("clCreateKernel") CHECK_ERROR("clCreateKernel")
pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
//parameters declaration //parameters declaration
int len; int len;
int depth; int depth;
...@@ -130,9 +131,9 @@ int main(int argc, char** argv) { ...@@ -130,9 +131,9 @@ int main(int argc, char** argv) {
h_Ax_vector=(float*)malloc(sizeof(float)*dim); h_Ax_vector=(float*)malloc(sizeof(float)*dim);
h_x_vector=(float*)malloc(sizeof(float)*dim); h_x_vector=(float*)malloc(sizeof(float)*dim);
input_vec( parameters->inpFiles[1],h_x_vector,dim); input_vec( parameters->inpFiles[1],h_x_vector,dim);
pb_SwitchToTimer(&timers, visc_TimerID_SETUP);
pb_SwitchToTimer(&timers, visc_TimerID_SETUP);
OpenCLDeviceProp clDeviceProp; OpenCLDeviceProp clDeviceProp;
clStatus = clGetDeviceInfo(clDevice,CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV,sizeof(cl_uint),&(clDeviceProp.major),NULL); clStatus = clGetDeviceInfo(clDevice,CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV,sizeof(cl_uint),&(clDeviceProp.major),NULL);
...@@ -206,14 +207,16 @@ pb_SwitchToTimer(&timers, pb_TimerID_COPY); ...@@ -206,14 +207,16 @@ pb_SwitchToTimer(&timers, pb_TimerID_COPY);
CHECK_ERROR("clSetKernelArg") CHECK_ERROR("clSetKernelArg")
//main execution //main execution
pb_SwitchToTimer(&timers, pb_TimerID_KERNEL); pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
printf("grid = %lu, block = %lu\n", grid, block); printf("grid = %lu, block = %lu\n", grid, block);
int i; int i;
for(i=0; i<50; i++) for(i=0; i<50; i++)
{ {
pb_SwitchToTimer(&timers, pb_TimerID_KERNEL);
clStatus = clEnqueueNDRangeKernel(clCommandQueue,clKernel,1,NULL,&grid,&block,0,NULL,NULL); clStatus = clEnqueueNDRangeKernel(clCommandQueue,clKernel,1,NULL,&grid,&block,0,NULL,NULL);
CHECK_ERROR("clEnqueueNDRangeKernel") CHECK_ERROR("clEnqueueNDRangeKernel")
pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
} }
clStatus = clFinish(clCommandQueue); clStatus = clFinish(clCommandQueue);
...@@ -226,6 +229,7 @@ pb_SwitchToTimer(&timers, pb_TimerID_COPY); ...@@ -226,6 +229,7 @@ pb_SwitchToTimer(&timers, pb_TimerID_COPY);
clStatus = clEnqueueReadBuffer(clCommandQueue,d_Ax_vector,CL_TRUE,0,dim*sizeof(float),h_Ax_vector,0,NULL,NULL); clStatus = clEnqueueReadBuffer(clCommandQueue,d_Ax_vector,CL_TRUE,0,dim*sizeof(float),h_Ax_vector,0,NULL,NULL);
CHECK_ERROR("clEnqueueReadBuffer") CHECK_ERROR("clEnqueueReadBuffer")
pb_SwitchToTimer(&timers, visc_TimerID_SETUP);
clStatus = clReleaseKernel(clKernel); clStatus = clReleaseKernel(clKernel);
clStatus = clReleaseProgram(clProgram); clStatus = clReleaseProgram(clProgram);
......
...@@ -158,16 +158,14 @@ int main(int argc, char** argv) { ...@@ -158,16 +158,14 @@ int main(int argc, char** argv) {
h_x_vector=(float*)malloc(sizeof(float)*dim); h_x_vector=(float*)malloc(sizeof(float)*dim);
input_vec( parameters->inpFiles[1],h_x_vector,dim); input_vec( parameters->inpFiles[1],h_x_vector,dim);
pb_SwitchToTimer(&timers, pb_TimerID_COPY); pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
memset(h_Ax_vector, 0, dim*sizeof(float)); memset(h_Ax_vector, 0, dim*sizeof(float));
size_t grid; size_t grid;
size_t block; size_t block;
pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
compute_active_thread(&block, &grid, nzcnt_len, pad, 3, 0, 8); compute_active_thread(&block, &grid, nzcnt_len, pad, 3, 0, 8);
pb_SwitchToTimer(&timers, visc_TimerID_MEM_TRACK); pb_SwitchToTimer(&timers, visc_TimerID_MEM_TRACK);
llvm_visc_track_mem(h_Ax_vector, dim*sizeof(float)); llvm_visc_track_mem(h_Ax_vector, dim*sizeof(float));
llvm_visc_track_mem(h_data, len*sizeof(float)); llvm_visc_track_mem(h_data, len*sizeof(float));
...@@ -180,12 +178,12 @@ int main(int argc, char** argv) { ...@@ -180,12 +178,12 @@ int main(int argc, char** argv) {
//main execution //main execution
pb_SwitchToTimer(&timers, pb_TimerID_KERNEL); pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
int i; int i;
for(i=0; i<50; i++) for(i=0; i<50; i++)
{ {
pb_SwitchToTimer(&timers, pb_TimerID_NONE);
unsigned spmvDFG = __visc__node(spmv_jds, 2, 1, (unsigned)block, (unsigned)(grid/block), unsigned spmvDFG = __visc__node(spmv_jds, 2, 1, (unsigned)block, (unsigned)(grid/block),
15, 15,
h_Ax_vector, h_Ax_vector,
...@@ -205,6 +203,7 @@ int main(int argc, char** argv) { ...@@ -205,6 +203,7 @@ int main(int argc, char** argv) {
nzcnt_len*sizeof(int), nzcnt_len*sizeof(int),
0); 0);
__visc__wait(spmvDFG); __visc__wait(spmvDFG);
pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
/******************************* Issues ******************************* /******************************* Issues *******************************
* 1. Using OpenCL to compute grid and block dimensions * 1. Using OpenCL to compute grid and block dimensions
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment