I kinda suspect he might be measuring the time it takes to launch a kernel rather than the time it takes the kernel to complete.
Thrust device calls, like those of the underlying CUDA library, are asynchronous by default. The only exception is calls that result in a memcpy, which are synchronous. To wait until an async call is completed you need to call one of the synchronize commands, like cudaDeviceSynchronize.
Looking through his test.cu file, he snaps a timestamp using std::clock right after doing the kernel launch with for_each. Ignoring the fact that this is not an accurate way to benchmark a GPU (you need to use events to accurately benchmark the kernel) what you're capturing will just be the processor time it takes to make the async kernel launch. Std::clock measures CPU time, which is (rightly) close to 0 for a program that runs on the GPU.
It's entirely possible that you're not even getting valid results out of the other end - note that you don't show output. I don't know if thrust's magic device memory access function triggers a synchronization or not. I kinda remember having to make an explicit call when I did a GPU simulation.
I don't have access to a CUDA box at the moment, I'd have to add those cudaDeviceSynchronize calls after the for_each invocations to be sure.
Thrust CUDA calls are synchronous to each other. You can add an explicit synchronization call cudaDeviceSynchronize() and there won't be a difference in results.
Comments
I kinda suspect he might be measuring the time it takes to launch a kernel rather than the time it takes the kernel to complete.
Thrust device calls, like those of the underlying CUDA library, are asynchronous by default. The only exception is calls that result in a memcpy, which are synchronous. To wait until an async call is completed you need to call one of the synchronize commands, like cudaDeviceSynchronize.
Looking through his test.cu file, he snaps a timestamp using std::clock right after doing the kernel launch with for_each. Ignoring the fact that this is not an accurate way to benchmark a GPU (you need to use events to accurately benchmark the kernel) what you're capturing will just be the processor time it takes to make the async kernel launch. Std::clock measures CPU time, which is (rightly) close to 0 for a program that runs on the GPU.
It's entirely possible that you're not even getting valid results out of the other end - note that you don't show output. I don't know if thrust's magic device memory access function triggers a synchronization or not. I kinda remember having to make an explicit call when I did a GPU simulation.
I don't have access to a CUDA box at the moment, I'd have to add those cudaDeviceSynchronize calls after the for_each invocations to be sure.
Thrust CUDA calls are synchronous to each other. You can add an explicit synchronization call cudaDeviceSynchronize() and there won't be a difference in results.
I had suspected this as well, but I'm not familiar with how Thrust's CUDA bindings work.