EDIT: So, I forgot to attach a picture, and now, for some reason I reloaded the data into excel and the peak no longer shifts from regular math to fast math…
Hmmmm
But, here is the picture now
Hi all,
I am trying to interpret the speedups I obtained as a function of the number of blocks sent to the GPU.
The occupancy calc says I have 8 blocks per MP, with 64 threads per block.
If I increase the number of blocks I thought I would have an increasing level of performance vs the serial, one-threaded CPU code which scales linearly with increasing elements to calculate (the same ‘elements’ that lead to the threads used in the block, so its a one-for-one comparison).
Instead I get the attached figure.
Please note that the option with fast-math peaks at 768 blocks, while the other is at 384 blocks.
Does anyone have an idea (or suggestions for what I should look further in to) why there is that peak at 384/768 blocks?
The fast math option results in no local mem being needed, the non-fast math one uses 168 bytes of lmem, if that matters.
The data points with < 384 points were run with numbers of elements such that the last block is partially not-filled, and since my kernel has something like "if (threadIdx.x+blockDim.x*blockIdx.x) <total_number_elements), there will be divergence added. At first this was partially the culprit, until I saw the fast math result.
Both codes are compiled with -arch sm_13
I am running on a GTX 275 (30 MPs).
Thanks very much,
Adam
