sequential sum within a kernel.

[E.D_Riedijk]

This is very interesting, but please, could you give me a few more hints?

3 questions:

1. Why my reads&writes are uncoalesced? I know they are, it’s just that I can’t figure out why in this case (a reference to the CUDA programming guide is OK, I did read it but may have missed something)

To coalesce, thread 0 of a half-warp needs to read memory location [N16+0], thread 1, needs to read [N16+1], etc. You are reading memory locations SegmentLen apart between the threads.

2. In my case ArraySize goes from 512512=262’144 to 20482048=4’194’304.

Since BLOCKSIZE can’t be more than 512, can I really set GRIDSIZE up to 8192?

Isn’t it too much?

With 16 SMs (G92 GTS), it would mean 512 blocks per SM (I was under the impression that we can’t have more than 24 blocks per SM?).

If you read the programming manual (or check devicequery), you can launch a grid of size 65535 x 65535. So yes, you can easily launch such a grid.

You cannot have more than x blocks at the same time on 1 multiprocessor (depends on block size and amount of registers used & shared memory). Other blocks get scheduled as soon as a block is finished. To get optimal performance out of CUDA, you normally want to launch thousands of blocks (or at least as much as possible)

3. Is there a way to exploit those 16KB of shared memory per SM, in my case? I don’t think so since it’s a per-element operation, but…

[snapback]436960[/snapback]

Indeed, for you there is no use.

[XFer]

Thank you for the quick reply. :)

To coalesce, thread 0 of a half-warp needs to read memory location [N16+0], thread 1, needs to read [N16+1], etc. You are reading memory locations SegmentLen apart between the threads.

In my 2048x2048 case, SegmentLen is 32 = 16*2; yet the profiler tells me I have fully uncoalesced reads and writes, and indeed the relative performances are no better than the 1024x1024 case (4x worse in fact, which is linear worsening).

Anyway this evening I’ll try the new code, thanks! External Media

Fernando

[E.D_Riedijk]

Thank you for the quick reply. :)

In my 2048x2048 case, SegmentLen is 32 = 16*2; yet the profiler tells me I have fully uncoalesced reads and writes, and indeed the relative performances are no better than the 1024x1024 case (4x worse in fact, which is linear worsening).

Anyway this evening I’ll try the new code, thanks!  External Media

Fernando

[snapback]436990[/snapback]

you will only coalesce if SegmentLen = 1 on anything older than GTX260/280. On GTX260/280, it will be approximately like this:

segmentlen = 2 → 2 times as slow

until segmentlen = 16. Then after that segmentlen = 32 and higher should have the same (slow) speed as segmentlen = 16.

[XFer]

Great! 10x the speed and Occupancy = 1!! External Image

Thanks a lot Denis!

For reference, I ended up with:

__global__ void  ArraySumInt (int *ResArray, const int *Array1, const int *Array2, const unsigned int iSize)

{

	

  unsigned int	index = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;

 if (index < iSize)

   ResArray[index] = Array1[index] + Array2[index];

}

Fernando