Poor heterogeneous memory performance in Quartus 14.1

You need two things to avoid having it allocated in "default" global memory. If you don't do these two things then you pay a penalty of the software reading it from DDR and writing it to QDR prior to kernel execution (depending on when you take your timestamp, this can look like lower kernel performance - check the profiler to be sure). But in short you need:

1. Use the cl_mem flag CL_MEM_HETEROGENEOUS_ALTERA when you do CreateBuffer

2. Use clSetKernelArg to bind it to the argument that has the "buffer_location" attribute BEFORE doing any accesses to that buffer.

For example you need to do:

mem = clCreateBuffer(context, flags|CL_MEM_HETEROGENEOUS_ALTERA, memSize, NULL,&errNum);

clSetKernelArg(kernel, 0, sizeof(cl_mem), &mem);

clEnqueueWriteBuffer(queue, mem, CL_FALSE, 0, N, 0, NULL, &write_event);

clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, NULL, 0, NULL, &kernel_event);

That should launch your kernel without any unnecessary copies.